Merge branch 'main' into cuda-arch-fixup

Signed-off-by: Shengqi Chen <harry-chen@outlook.com>

CMakeLists.txt

@@ -227,6 +227,22 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   cuda_archs_loose_intersection(CUDA_ARCHS
     "${CUDA_SUPPORTED_ARCHS}" "${CUDA_ARCHS}")
   message(STATUS "CUDA supported target architectures: ${CUDA_ARCHS}")
+
+  set(VLLM_COMPILED_CUDA_ARCHS)
+  foreach(_ARCH ${CUDA_ARCHS})
+    set(_COMPILED_ARCH "${_ARCH}")
+    if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
+      if(_ARCH MATCHES "^(10|11|12)\\.0$")
+        set(_COMPILED_ARCH "${_ARCH}f")
+      endif()
+    elseif(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.8)
+      if(_ARCH MATCHES "^(10\\.(0|1|3)|12\\.(0|1))$")
+        set(_COMPILED_ARCH "${_ARCH}a")
+      endif()
+    endif()
+    list(APPEND VLLM_COMPILED_CUDA_ARCHS "${_COMPILED_ARCH}")
+  endforeach()
+  list(JOIN VLLM_COMPILED_CUDA_ARCHS "," VLLM_COMPILED_CUDA_ARCHS_STR)
 else()
   #
   # For other GPU targets override the GPU architectures detected by cmake/torch
@@ -385,8 +401,20 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   #
   # _C_stable_libtorch extension (ops registered via STABLE_TORCH_LIBRARY)
   #
+  # Shared entry sources are part of the base extension source list, but some
+  # optional kernel families below append source-local feature macros to them.
+  set(VLLM_STABLE_TORCH_BINDINGS_SRC
+    "csrc/libtorch_stable/torch_bindings.cpp")
+  set(CACHE_KERNELS_SRC "csrc/libtorch_stable/cache_kernels.cu")
+  set(SCALED_MM_ENTRY_SRC
+    "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu")
+  set(NVFP4_QUANT_ENTRY_SRC
+    "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu")
+  set(NVFP4_SCALED_MM_ENTRY_SRC
+    "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_entry.cu")
+
   set(VLLM_STABLE_EXT_SRC
-    "csrc/libtorch_stable/torch_bindings.cpp"
+    "${VLLM_STABLE_TORCH_BINDINGS_SRC}"
     "csrc/libtorch_stable/cuda_view.cu"
     "csrc/libtorch_stable/cuda_utils_kernels.cu"
     "csrc/libtorch_stable/activation_kernels.cu"
@@ -409,7 +437,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     "csrc/libtorch_stable/sampler.cu"
     "csrc/libtorch_stable/topk.cu"
     "csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
-    "csrc/libtorch_stable/cache_kernels.cu"
+    "${CACHE_KERNELS_SRC}"
     "csrc/libtorch_stable/cache_kernels_fused.cu"
     "csrc/libtorch_stable/custom_all_reduce.cu"
     "csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
@@ -425,11 +453,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       cuda_archs_loose_intersection(COOPERATIVE_TOPK_ARCHS
         "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
     endif()
-
-    if(COOPERATIVE_TOPK_ARCHS)
-      list(APPEND VLLM_GPU_FLAGS "-DVLLM_ENABLE_COOPERATIVE_TOPK=1")
-
-    endif()
   endif()
 
   if(VLLM_GPU_LANG STREQUAL "CUDA")
@@ -467,11 +490,14 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
 
     list(APPEND VLLM_STABLE_EXT_SRC
       "csrc/libtorch_stable/cutlass_extensions/common.cpp"
-      "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_entry.cu"
+      "${SCALED_MM_ENTRY_SRC}"
+      "${NVFP4_QUANT_ENTRY_SRC}"
+      "${NVFP4_SCALED_MM_ENTRY_SRC}"
       "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu"
       "csrc/libtorch_stable/minimax_reduce_rms_kernel.cu")
+    set_compile_definitions_for_srcs(
+      SRCS "${VLLM_STABLE_TORCH_BINDINGS_SRC}"
+      DEFINITIONS VLLM_COMPILED_CUDA_ARCHS=\"${VLLM_COMPILED_CUDA_ARCHS_STR}\")
 
     #
     # Machete kernels
@@ -547,11 +573,15 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       CUDA_ARCHS "${CUDA_ARCHS}")
 
     if(COOPERATIVE_TOPK_ARCHS)
+      set(COOPERATIVE_TOPK_SRC "csrc/libtorch_stable/cooperative_topk.cu")
       list(APPEND VLLM_STABLE_EXT_SRC
-        "csrc/libtorch_stable/cooperative_topk.cu")
+        "${COOPERATIVE_TOPK_SRC}")
       set_gencode_flags_for_srcs(
-        SRCS "csrc/libtorch_stable/cooperative_topk.cu"
+        SRCS "${COOPERATIVE_TOPK_SRC}"
         CUDA_ARCHS "${COOPERATIVE_TOPK_ARCHS}")
+      set_compile_definitions_for_srcs(
+        SRCS "${VLLM_STABLE_TORCH_BINDINGS_SRC};${COOPERATIVE_TOPK_SRC}"
+        DEFINITIONS VLLM_ENABLE_COOPERATIVE_TOPK=1)
     endif()
 
     # Only build Marlin kernels if we are building for at least some compatible archs.
@@ -760,8 +790,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_SM90_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${SCALED_MM_ENTRY_SRC};${SCALED_MM_SM90_SRCS}"
+      DEFINITIONS ENABLE_SCALED_MM_SM90=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_SM90_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_SM90=1")
     # Let scaled_mm_c2x know it doesn't need to build these arches
     list(APPEND SCALED_MM_3X_ARCHS "${SCALED_MM_ARCHS}")
     message(STATUS "Building scaled_mm_c3x_sm90 for archs: ${SCALED_MM_ARCHS}")
@@ -794,8 +826,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_SM120_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${SCALED_MM_ENTRY_SRC};${SCALED_MM_SM120_SRCS}"
+      DEFINITIONS ENABLE_SCALED_MM_SM120=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_SM120_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_SM120=1")
     # Let scaled_mm_c2x know it doesn't need to build these arches
     list(APPEND SCALED_MM_3X_ARCHS "${SCALED_MM_ARCHS}")
     message(STATUS "Building scaled_mm_c3x_sm120 for archs: ${SCALED_MM_ARCHS}")
@@ -828,8 +862,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_SM100_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${SCALED_MM_ENTRY_SRC};${SCALED_MM_SM100_SRCS}"
+      DEFINITIONS ENABLE_SCALED_MM_SM100=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_SM100_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_SM100=1")
     # Let scaled_mm_c2x know it doesn't need to build these arches
     list(APPEND SCALED_MM_3X_ARCHS "${SCALED_MM_ARCHS}")
     message(STATUS "Building scaled_mm_c3x_sm100 for archs: ${SCALED_MM_ARCHS}")
@@ -858,8 +894,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set_gencode_flags_for_srcs(
       SRCS "${SCALED_MM_C2X_SRCS}"
       CUDA_ARCHS "${SCALED_MM_2X_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${SCALED_MM_ENTRY_SRC};${SCALED_MM_C2X_SRCS}"
+      DEFINITIONS ENABLE_SCALED_MM_C2X=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_C2X_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_C2X=1")
     message(STATUS "Building scaled_mm_c2x for archs: ${SCALED_MM_2X_ARCHS}")
   else()
     if (SCALED_MM_3X_ARCHS)
@@ -884,8 +922,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set_gencode_flags_for_srcs(
       SRCS "${CUTLASS_MOE_SM90_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${SCALED_MM_ENTRY_SRC};${CUTLASS_MOE_SM90_SRCS}"
+      DEFINITIONS ENABLE_CUTLASS_MOE_SM90=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${CUTLASS_MOE_SM90_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM90=1")
     message(STATUS "Building grouped_mm_c3x for archs: ${SCALED_MM_ARCHS}")
   else()
     if (NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.3 AND SCALED_MM_ARCHS)
@@ -908,8 +948,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set_gencode_flags_for_srcs(
       SRCS "${CUTLASS_MOE_SM100_SRCS}"
       CUDA_ARCHS "${SCALED_MM_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${SCALED_MM_ENTRY_SRC};${CUTLASS_MOE_SM100_SRCS}"
+      DEFINITIONS ENABLE_CUTLASS_MOE_SM10X_OR_SM11X=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${CUTLASS_MOE_SM100_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM100=1")
+    # The implementation is named sm100 historically, but it is built for the
+    # SM10x/SM11x family: CUDA 12 Thor reports SM101, CUDA 13 Thor reports
+    # SM110. Keep the compile-time macro aligned with runtime dispatch.
     message(STATUS "Building grouped_mm_c3x for archs: ${SCALED_MM_ARCHS}")
   else()
     if (NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
@@ -950,10 +995,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # FP4/NVFP4 kernels (moved from _C to _C_stable_libtorch)
   #
 
-  # SM12x FP4 kernels. These share some generic NVFP4 quantization entry
-  # sources with the SM10x/11x block below; set_gencode_flags_for_srcs appends
-  # per-source flags, so shared files accumulate both SM12x and SM10x/11x
-  # gencodes when both families are requested.
+  # Shared FP4 implementation sources live next to the FP4 arch logic because
+  # both SM12x and SM10x/11x append family-specific gencodes and feature macros
+  # to them.
+  set(FP4_SHARED_SRCS
+    "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu"
+    "csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
+    "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
+    "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
+    "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
+
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
     cuda_archs_loose_intersection(FP4_SM120_ARCHS "12.0f" "${CUDA_ARCHS}")
   else()
@@ -961,18 +1012,19 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_SM120_ARCHS)
     set(FP4_SM120_SRCS
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
+      ${FP4_SHARED_SRCS}
       "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
+    )
     set_gencode_flags_for_srcs(
       SRCS "${FP4_SM120_SRCS}"
       CUDA_ARCHS "${FP4_SM120_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${NVFP4_QUANT_ENTRY_SRC};${NVFP4_SCALED_MM_ENTRY_SRC};${CACHE_KERNELS_SRC};${FP4_SM120_SRCS}"
+      DEFINITIONS ENABLE_NVFP4_SM120=1)
+    set_compile_definitions_for_srcs(
+      SRCS "${SCALED_MM_ENTRY_SRC}"
+      DEFINITIONS ENABLE_CUTLASS_MOE_SM120=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${FP4_SM120_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM120=1")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM120=1")
     message(STATUS "Building SM12x NVFP4 for archs: ${FP4_SM120_ARCHS}")
   else()
     message(STATUS "Not building SM12x NVFP4 as no compatible archs were found.")
@@ -987,14 +1039,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_SM100_ARCHS)
     set(FP4_SM100_SRCS
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
+      ${FP4_SHARED_SRCS}
       "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
-      "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
       "csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu"
-      "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu"
-      "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
+      "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu")
     if(NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9)
       message(STATUS
         "Building mxfp4_experts_quant unsupported stubs because CUDA compiler version is not >= 12.9 (found ${CMAKE_CUDA_COMPILER_VERSION}).")
@@ -1002,9 +1050,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set_gencode_flags_for_srcs(
       SRCS "${FP4_SM100_SRCS}"
       CUDA_ARCHS "${FP4_SM100_ARCHS}")
+    set_compile_definitions_for_srcs(
+      SRCS "${NVFP4_QUANT_ENTRY_SRC};${NVFP4_SCALED_MM_ENTRY_SRC};${CACHE_KERNELS_SRC};${FP4_SM100_SRCS}"
+      DEFINITIONS ENABLE_NVFP4_SM100=1)
     list(APPEND VLLM_STABLE_EXT_SRC "${FP4_SM100_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM100=1")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM100=1")
     message(STATUS "Building SM10x/11x NVFP4/MXFP4 for archs: ${FP4_SM100_ARCHS}")
   else()
     message(STATUS "Not building SM10x/11x NVFP4/MXFP4 as no compatible archs were found.")
@@ -1058,7 +1107,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
       SRCS "${CUTLASS_MLA_SRCS}"
       CUDA_ARCHS "${MLA_ARCHS}")
     list(APPEND VLLM_STABLE_EXT_SRC "${CUTLASS_MLA_SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
     # Add MLA-specific include directories only to MLA source files
     set_source_files_properties(${CUTLASS_MLA_SRCS}
       PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
@@ -1106,10 +1154,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # Needed to use cuda/hip APIs from C-shim
   if(VLLM_GPU_LANG STREQUAL "CUDA")
     target_compile_definitions(_C_stable_libtorch PRIVATE USE_CUDA)
-    if(COOPERATIVE_TOPK_ARCHS)
-      target_compile_definitions(_C_stable_libtorch PRIVATE
-        VLLM_ENABLE_COOPERATIVE_TOPK=1)
-    endif()
     # Needed by CUTLASS kernels
     target_compile_definitions(_C_stable_libtorch PRIVATE
       CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)

cmake/utils.cmake

@@ -344,6 +344,28 @@ macro(set_gencode_flags_for_srcs)
   endif()
 endmacro()
 
+#
+# For a list of source files append preprocessor definitions to file-specific
+# compile options. Use this for optional kernel feature macros so toggling one
+# kernel family does not perturb the compile command for every source in the
+# extension target.
+#
+macro(set_compile_definitions_for_srcs)
+  set(options)
+  set(oneValueArgs)
+  set(multiValueArgs SRCS DEFINITIONS)
+  cmake_parse_arguments(arg "${options}" "${oneValueArgs}"
+                        "${multiValueArgs}" ${ARGN})
+
+  foreach(_DEF ${arg_DEFINITIONS})
+    set_property(
+      SOURCE ${arg_SRCS}
+      APPEND PROPERTY
+      COMPILE_DEFINITIONS "${_DEF}"
+    )
+  endforeach()
+endmacro()
+
 #
 # For the given `SRC_CUDA_ARCHS` list of gencode versions in the form
 #  `<major>.<minor>[letter]` compute the "loose intersection" with the

csrc/libtorch_stable/ops.h

@@ -47,6 +47,8 @@ torch::stable::Tensor permute_cols(torch::stable::Tensor const& A,
                                    torch::stable::Tensor const& perm);
 
 #ifndef USE_ROCM
+std::string get_compiled_cuda_archs();
+
 bool cutlass_scaled_mm_supports_fp8(int64_t cuda_device_capability);
 bool cutlass_scaled_mm_supports_block_fp8(int64_t cuda_device_capability);
 bool cutlass_group_gemm_supported(int64_t cuda_device_capability);

csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu

@@ -51,7 +51,8 @@ void cutlass_moe_mm_sm90(torch::stable::Tensor& out_tensors,
 
 #endif
 
-#if defined ENABLE_CUTLASS_MOE_SM100 && ENABLE_CUTLASS_MOE_SM100
+#if defined ENABLE_CUTLASS_MOE_SM10X_OR_SM11X && \
+    ENABLE_CUTLASS_MOE_SM10X_OR_SM11X
 void cutlass_moe_mm_sm100(torch::stable::Tensor& out_tensors,
                           torch::stable::Tensor const& a_tensors,
                           torch::stable::Tensor const& b_tensors,
@@ -83,8 +84,9 @@ void cutlass_scaled_mm_sm100(torch::stable::Tensor& c,
                              std::optional<torch::stable::Tensor> const& bias);
 #endif
 
-#if (defined(ENABLE_CUTLASS_MOE_SM90) && ENABLE_CUTLASS_MOE_SM90) ||   \
-    (defined(ENABLE_CUTLASS_MOE_SM100) && ENABLE_CUTLASS_MOE_SM100) || \
+#if (defined(ENABLE_CUTLASS_MOE_SM90) && ENABLE_CUTLASS_MOE_SM90) || \
+    (defined(ENABLE_CUTLASS_MOE_SM10X_OR_SM11X) &&                   \
+     ENABLE_CUTLASS_MOE_SM10X_OR_SM11X) ||                           \
     (defined(ENABLE_CUTLASS_MOE_SM120) && ENABLE_CUTLASS_MOE_SM120)
 void get_cutlass_moe_mm_data_caller(
     const torch::stable::Tensor& topk_ids,
@@ -175,11 +177,14 @@ bool cutlass_scaled_mm_supports_block_fp8(int64_t cuda_device_capability) {
 
 bool cutlass_group_gemm_supported(int64_t cuda_device_capability) {
   // CUTLASS grouped FP8 kernels need at least CUDA 12.3 and SM90 (Hopper)
-  // or CUDA 12.8 and SM100 (Blackwell). Only report archs that have an
-  // actual cutlass_moe_mm dispatch compiled into this file.
+  // or CUDA 12.8 and SM10x/SM11x (Blackwell / Thor). CUDA 12 reports Thor as
+  // SM101 while CUDA 13 reports it as SM110, but both use this sm100-named
+  // implementation. Only report archs that have an actual cutlass_moe_mm
+  // dispatch compiled into this file.
 
 #if defined CUDA_VERSION
-  #if defined ENABLE_CUTLASS_MOE_SM100 && ENABLE_CUTLASS_MOE_SM100
+  #if defined ENABLE_CUTLASS_MOE_SM10X_OR_SM11X && \
+      ENABLE_CUTLASS_MOE_SM10X_OR_SM11X
   if (cuda_device_capability >= 100 && cuda_device_capability < 120) {
     return CUDA_VERSION >= 12080;
   }
@@ -281,8 +286,11 @@ void cutlass_moe_mm(torch::stable::Tensor& out_tensors,
                     torch::stable::Tensor const& c_strides, bool per_act_token,
                     bool per_out_ch) {
   int32_t version_num = get_sm_version_num();
-#if defined ENABLE_CUTLASS_MOE_SM100 && ENABLE_CUTLASS_MOE_SM100
-  if (version_num >= 100 && version_num < 110) {
+#if defined ENABLE_CUTLASS_MOE_SM10X_OR_SM11X && \
+    ENABLE_CUTLASS_MOE_SM10X_OR_SM11X
+  // Keep runtime dispatch aligned with the CMake arch list and support query:
+  // CUDA 12 Thor is SM101 and CUDA 13 Thor is SM110.
+  if (version_num >= 100 && version_num < 120) {
     cutlass_moe_mm_sm100(out_tensors, a_tensors, b_tensors, a_scales, b_scales,
                          expert_offsets, problem_sizes, a_strides, b_strides,
                          c_strides, per_act_token, per_out_ch);
@@ -316,8 +324,9 @@ void get_cutlass_moe_mm_data(
   // This function currently gets compiled only if we have a valid cutlass moe
   // mm to run it for.
   int32_t version_num = get_sm_version_num();
-#if (defined ENABLE_CUTLASS_MOE_SM90 && ENABLE_CUTLASS_MOE_SM90) ||   \
-    (defined ENABLE_CUTLASS_MOE_SM100 && ENABLE_CUTLASS_MOE_SM100) || \
+#if (defined ENABLE_CUTLASS_MOE_SM90 && ENABLE_CUTLASS_MOE_SM90) || \
+    (defined ENABLE_CUTLASS_MOE_SM10X_OR_SM11X &&                   \
+     ENABLE_CUTLASS_MOE_SM10X_OR_SM11X) ||                          \
     (defined ENABLE_CUTLASS_MOE_SM120 && ENABLE_CUTLASS_MOE_SM120)
   get_cutlass_moe_mm_data_caller(topk_ids, expert_offsets, problem_sizes1,
                                  problem_sizes2, input_permutation,
@@ -338,8 +347,9 @@ void get_cutlass_moe_mm_problem_sizes_from_expert_offsets(
     torch::stable::Tensor& problem_sizes2, const int64_t n, const int64_t k,
     const bool swap_ab) {
   int32_t version_num = get_sm_version_num();
-#if (defined ENABLE_CUTLASS_MOE_SM90 && ENABLE_CUTLASS_MOE_SM90) ||   \
-    (defined ENABLE_CUTLASS_MOE_SM100 && ENABLE_CUTLASS_MOE_SM100) || \
+#if (defined ENABLE_CUTLASS_MOE_SM90 && ENABLE_CUTLASS_MOE_SM90) || \
+    (defined ENABLE_CUTLASS_MOE_SM10X_OR_SM11X &&                   \
+     ENABLE_CUTLASS_MOE_SM10X_OR_SM11X) ||                          \
     (defined ENABLE_CUTLASS_MOE_SM120 && ENABLE_CUTLASS_MOE_SM120)
   get_cutlass_moe_mm_problem_sizes_from_expert_offsets_caller(
       expert_first_token_offset, problem_sizes1, problem_sizes2, n, k, swap_ab);
@@ -362,8 +372,9 @@ void get_cutlass_batched_moe_mm_data(
   // This function currently gets compiled only if we have a valid cutlass moe
   // mm to run it for.
   int32_t version_num = get_sm_version_num();
-#if (defined ENABLE_CUTLASS_MOE_SM90 && ENABLE_CUTLASS_MOE_SM90) ||   \
-    (defined ENABLE_CUTLASS_MOE_SM100 && ENABLE_CUTLASS_MOE_SM100) || \
+#if (defined ENABLE_CUTLASS_MOE_SM90 && ENABLE_CUTLASS_MOE_SM90) || \
+    (defined ENABLE_CUTLASS_MOE_SM10X_OR_SM11X &&                   \
+     ENABLE_CUTLASS_MOE_SM10X_OR_SM11X) ||                          \
     (defined ENABLE_CUTLASS_MOE_SM120 && ENABLE_CUTLASS_MOE_SM120)
   get_cutlass_batched_moe_mm_data_caller(expert_offsets, problem_sizes1,
                                          problem_sizes2, expert_num_tokens,

csrc/libtorch_stable/torch_bindings.cpp

@@ -4,6 +4,10 @@
 
 #include <torch/csrc/stable/library.h>
 
+#ifndef USE_ROCM
+std::string get_compiled_cuda_archs() { return VLLM_COMPILED_CUDA_ARCHS; }
+#endif
+
 // Register ops with STABLE_TORCH_LIBRARY for libtorch stable ABI compatibility.
 // Note: We register under namespace "_C" so ops are accessible as
 // torch.ops._C.<op_name> for compatibility with existing code.
@@ -30,6 +34,7 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
   ops.def("permute_cols(Tensor A, Tensor perm) -> Tensor");
 
   ops.def("get_cuda_view_from_cpu_tensor(Tensor cpu_tensor) -> Tensor");
+  ops.def("get_compiled_cuda_archs() -> str");
 
 #ifndef USE_ROCM
 
@@ -763,6 +768,7 @@ STABLE_TORCH_LIBRARY_IMPL(_C_cuda_utils, CompositeExplicitAutograd,
 // ops.impl("op_name", &func) without a dispatch key in the non-stable API.
 STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, ops) {
 #ifndef USE_ROCM
+  ops.impl("get_compiled_cuda_archs", TORCH_BOX(&get_compiled_cuda_archs));
   ops.impl("cutlass_scaled_mm_supports_fp8",
            TORCH_BOX(&cutlass_scaled_mm_supports_fp8));
   ops.impl("cutlass_group_gemm_supported",

docs/getting_started/installation/gpu.cuda.inc.md

@@ -46,6 +46,26 @@ export CPU_ARCH=$(uname -m) # x86_64 or aarch64
 uv pip install https://github.com/vllm-project/vllm/releases/download/v${VLLM_VERSION}/vllm-${VLLM_VERSION}+cu${CUDA_VERSION}-cp38-abi3-manylinux_2_28_${CPU_ARCH}.whl --extra-index-url https://download.pytorch.org/whl/cu${CUDA_VERSION}
 ```
 
+!!! warning "CUDA architecture coverage"
+
+    Pre-built CUDA wheels are compiled for the CUDA architectures selected by
+    vLLM's release and build pipelines. This list is intentionally smaller than
+    every architecture CMake can build, because each additional architecture
+    increases wheel size.
+
+    In particular, CUDA 12.9 wheels do not use CUDA 13 family-specific targets.
+    To keep wheel size bounded, published CUDA 12.9 wheels may omit some newer
+    architecture-specific Blackwell/Thor targets, such as `sm_103` or
+    `sm_121`, even though vLLM can build them from source. CUDA 13 wheels use
+    family-specific targets such as `sm_100f`, `sm_110f`, and `sm_120f`, which
+    cover the corresponding major-version GPU family.
+
+    If vLLM logs a warning that your visible CUDA device is not covered by the
+    wheel's compiled CUDA architectures, or if you see a CUDA error such as
+    `no kernel image is available for execution on the device`, install a CUDA
+    13 wheel when possible, or build from source with a `TORCH_CUDA_ARCH_LIST`
+    that includes your GPU.
+
 #### Install the latest code
 
 LLM inference is a fast-evolving field, and the latest code may contain bug fixes, performance improvements, and new features that are not released yet. To allow users to try the latest code without waiting for the next release, vLLM provides wheels for every commit since `v0.5.3` on <https://wheels.vllm.ai/nightly>. There are multiple indices that could be used:

docs/usage/troubleshooting.md

@@ -174,6 +174,33 @@ If the script runs successfully, you should see the message `sanity check is suc
 
 If the test script hangs or crashes, usually it means the hardware/drivers are broken in some sense. You should try to contact your system administrator or hardware vendor for further assistance. As a common workaround, you can try to tune some NCCL environment variables, such as `export NCCL_P2P_DISABLE=1` to see if it helps. Please check [their documentation](https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/env.html) for more information. Please only use these environment variables as a temporary workaround, as they might affect the performance of the system. The best solution is still to fix the hardware/drivers so that the test script can run successfully.
 
+## CUDA architecture not covered by the wheel
+
+If vLLM logs a warning that the current wheel was built for a set of CUDA
+architectures but one of your visible CUDA devices is not covered, the installed
+wheel may not contain native CUDA kernels for that GPU. The same issue may also
+surface later as a CUDA runtime error such as `no kernel image is available for
+execution on the device`.
+
+The architecture list in a pre-built wheel is determined by the vLLM release and
+build pipelines, then filtered by CMake before individual kernels choose their
+own per-kernel architectures. It is not the same as the full set of
+architectures that vLLM can build from source.
+
+CUDA 12.9 wheels use architecture-specific targets for newer NVIDIA GPUs. To
+keep wheel size bounded, published CUDA 12.9 wheels may omit some newer
+Blackwell/Thor targets such as `sm_103` or `sm_121`, even though vLLM can build
+them from source. CUDA 13 wheels can use family-specific targets such as
+`sm_100f`, `sm_110f`, and `sm_120f`, which cover the corresponding
+major-version GPU family.
+
+If you see this warning or error, install a CUDA 13 wheel when possible.
+Otherwise, build vLLM from source and set `TORCH_CUDA_ARCH_LIST` to include
+your GPU's compute capability. See the CUDA installation guide's
+[pre-built wheels](../getting_started/installation/gpu.md#pre-built-wheels) and
+[build wheel from source](../getting_started/installation/gpu.md#build-wheel-from-source)
+sections for details.
+
 ## Python multiprocessing
 
 ### `RuntimeError` Exception

tests/cuda/test_cuda_arch_warnings.py

@@ -0,0 +1,131 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import sys
+from types import ModuleType, SimpleNamespace
+from typing import Any
+
+import pytest
+
+from vllm.platforms.interface import DeviceCapability
+
+
+@pytest.fixture
+def cuda_platform_base(monkeypatch: pytest.MonkeyPatch) -> Any:
+    stable_libtorch_module = ModuleType("vllm._C_stable_libtorch")
+    monkeypatch.setitem(
+        sys.modules,
+        "vllm._C_stable_libtorch",
+        stable_libtorch_module,
+    )
+    from vllm.platforms.cuda import CudaPlatformBase
+
+    return CudaPlatformBase
+
+
+def test_compiled_arch_covers_device(cuda_platform_base: Any) -> None:
+    assert cuda_platform_base._compiled_arch_covers_device(
+        "12.1a", DeviceCapability(12, 1)
+    )
+    assert not cuda_platform_base._compiled_arch_covers_device(
+        "12.0a", DeviceCapability(12, 1)
+    )
+
+    assert cuda_platform_base._compiled_arch_covers_device(
+        "12.0f", DeviceCapability(12, 1)
+    )
+    assert not cuda_platform_base._compiled_arch_covers_device(
+        "10.0f", DeviceCapability(12, 1)
+    )
+
+
+def test_warn_if_device_arch_not_compiled(
+    monkeypatch: pytest.MonkeyPatch, cuda_platform_base: Any
+) -> None:
+    def device_count(cls: type[Any]) -> int:
+        return 2
+
+    def get_device_capability(
+        cls: type[Any], device_id: int = 0
+    ) -> DeviceCapability | None:
+        capabilities = {
+            0: DeviceCapability(12, 1),
+            1: DeviceCapability(10, 3),
+        }
+        return capabilities[device_id]
+
+    def get_device_name(cls: type[Any], device_id: int = 0) -> str:
+        return f"GPU {device_id}"
+
+    warnings: list[tuple[str, str, str]] = []
+
+    def warning_once(message: str, compiled_archs: str, devices: str) -> None:
+        warnings.append((message, compiled_archs, devices))
+
+    monkeypatch.setattr(cuda_platform_base, "device_count", classmethod(device_count))
+    monkeypatch.setattr(
+        cuda_platform_base,
+        "get_device_capability",
+        classmethod(get_device_capability),
+    )
+    monkeypatch.setattr(
+        cuda_platform_base, "get_device_name", classmethod(get_device_name)
+    )
+    monkeypatch.setattr(
+        "vllm.platforms.cuda.torch.ops",
+        SimpleNamespace(
+            _C=SimpleNamespace(get_compiled_cuda_archs=lambda: "12.0f,10.0a")
+        ),
+    )
+    monkeypatch.setattr(
+        "vllm.platforms.cuda.logger",
+        SimpleNamespace(warning_once=warning_once),
+    )
+
+    cuda_platform_base._warn_if_device_arch_not_compiled()
+
+    assert len(warnings) == 1
+    assert warnings[0][1] == "12.0f, 10.0a"
+    assert "1: GPU 1 (compute capability 10.3)" in warnings[0][2]
+
+
+def test_warn_if_device_arch_not_compiled_no_warning(
+    monkeypatch: pytest.MonkeyPatch, cuda_platform_base: Any
+) -> None:
+    def device_count(cls: type[Any]) -> int:
+        return 1
+
+    def get_device_capability(
+        cls: type[Any], device_id: int = 0
+    ) -> DeviceCapability | None:
+        return DeviceCapability(10, 3)
+
+    def get_device_name(cls: type[Any], device_id: int = 0) -> str:
+        raise AssertionError("get_device_name should not be called for covered devices")
+
+    warnings: list[tuple[str, str, str]] = []
+
+    def warning_once(message: str, compiled_archs: str, devices: str) -> None:
+        warnings.append((message, compiled_archs, devices))
+
+    monkeypatch.setattr(cuda_platform_base, "device_count", classmethod(device_count))
+    monkeypatch.setattr(
+        cuda_platform_base,
+        "get_device_capability",
+        classmethod(get_device_capability),
+    )
+    monkeypatch.setattr(
+        cuda_platform_base, "get_device_name", classmethod(get_device_name)
+    )
+    monkeypatch.setattr(
+        "vllm.platforms.cuda.torch.ops",
+        SimpleNamespace(_C=SimpleNamespace(get_compiled_cuda_archs=lambda: "10.0f")),
+    )
+    monkeypatch.setattr(
+        "vllm.platforms.cuda.logger",
+        SimpleNamespace(warning_once=warning_once),
+    )
+
+    cuda_platform_base._warn_if_device_arch_not_compiled()
+
+    assert warnings == []

tests/kernels/quantization/test_cutlass_group_gemm_support.py

@@ -0,0 +1,31 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+from types import SimpleNamespace
+
+import torch
+
+from vllm import _custom_ops as ops
+
+
+def test_cutlass_group_gemm_python_guard_allows_thor(monkeypatch):
+    seen_capabilities: list[int] = []
+
+    def fake_cutlass_group_gemm_supported(capability: int) -> bool:
+        seen_capabilities.append(capability)
+        return True
+
+    monkeypatch.setattr(
+        torch.ops,
+        "_C",
+        SimpleNamespace(cutlass_group_gemm_supported=fake_cutlass_group_gemm_supported),
+    )
+
+    # CUDA 12 reports Thor as SM101 and CUDA 13 reports it as SM110. Both
+    # should reach the C++ query for the SM10x/SM11x CUTLASS MoE kernel.
+    assert ops.cutlass_group_gemm_supported(101)
+    assert ops.cutlass_group_gemm_supported(110)
+
+    # SM120 uses separate kernels and must not be advertised by this path.
+    assert not ops.cutlass_group_gemm_supported(120)
+    assert seen_capabilities == [101, 110]

vllm/_custom_ops.py

@@ -835,12 +835,16 @@ def cutlass_scaled_mm_azp(
 
 
 def cutlass_group_gemm_supported(cuda_device_capability: int) -> bool:
-    if cuda_device_capability < 90 or cuda_device_capability >= 110:
+    # CUTLASS grouped FP8 MoE uses the same sm100-named implementation for
+    # SM10x and SM11x. Thor reports SM101 with CUDA 12 and SM110 with CUDA 13,
+    # so keep this Python guard aligned with the C++ support query/dispatch.
+    if cuda_device_capability < 90 or cuda_device_capability >= 120:
         return False
     try:
         return torch.ops._C.cutlass_group_gemm_supported(cuda_device_capability)
-    except AttributeError:
-        # Return False on non-CUDA platforms where it is not available
+    except (AttributeError, RuntimeError):
+        # Return False on non-CUDA platforms where it is not available or not
+        # implemented for the current build.
         return False
 
 

vllm/platforms/cuda.py

@@ -14,6 +14,7 @@ from datetime import timedelta
 from functools import cache, lru_cache, wraps
 from typing import TYPE_CHECKING, NamedTuple, TypeVar
 
+import regex as re
 import torch
 from torch.distributed import PrefixStore, ProcessGroup
 from torch.distributed.distributed_c10d import is_nccl_available
@@ -48,6 +49,8 @@ _R = TypeVar("_R")
 
 pynvml = import_pynvml()
 
+_COMPILED_CUDA_ARCH_RE = re.compile(r"^(?P<major>\d+)\.(?P<minor>\d+)(?P<kind>[af])?$")
+
 # pytorch 2.5 uses cudnn sdpa by default, which will cause crash on some models
 # see https://github.com/huggingface/diffusers/issues/9704 for details
 torch.backends.cuda.enable_cudnn_sdp(False)
@@ -282,6 +285,62 @@ class CudaPlatformBase(Platform):
     def log_warnings(cls):
         pass
 
+    @staticmethod
+    def _compiled_arch_covers_device(arch: str, capability: DeviceCapability) -> bool:
+        match = _COMPILED_CUDA_ARCH_RE.match(arch)
+        if match is None:
+            return False
+
+        major = int(match.group("major"))
+        minor = int(match.group("minor"))
+        kind = match.group("kind")
+
+        if kind == "f":
+            return major == capability.major
+        return major == capability.major and minor == capability.minor
+
+    @classmethod
+    def _warn_if_device_arch_not_compiled(cls) -> None:
+        try:
+            compiled_archs_str = torch.ops._C.get_compiled_cuda_archs()
+        except (AttributeError, RuntimeError):
+            return
+
+        compiled_archs = [
+            arch.strip() for arch in compiled_archs_str.split(",") if arch.strip()
+        ]
+        if not compiled_archs:
+            return
+
+        unsupported_devices: list[str] = []
+        for device_id in range(cls.device_count()):
+            capability = cls.get_device_capability(device_id)
+            if capability is None:
+                continue
+            if any(
+                cls._compiled_arch_covers_device(arch, capability)
+                for arch in compiled_archs
+            ):
+                continue
+            device_name = cls.get_device_name(device_id)
+            unsupported_devices.append(
+                f"{device_id}: {device_name} (compute capability "
+                f"{capability.as_version_str()})"
+            )
+
+        if unsupported_devices:
+            logger.warning_once(
+                "The current vLLM wheel was built for CUDA architectures %s, "
+                "but the following visible CUDA device(s) are not covered: %s. "
+                "This may cause CUDA kernel launch failures or force slower "
+                "fallback paths. Note that CUDA 13 family-specific targets "
+                "such as 10.0f and 12.0f cover the corresponding major-version "
+                "GPU family, while architecture-specific targets such as 12.1a "
+                "cover only that exact compute capability.",
+                ", ".join(compiled_archs),
+                "; ".join(unsupported_devices),
+            )
+
     @classmethod
     def is_pin_memory_available(cls) -> bool:
         if in_wsl():
@@ -949,6 +1008,7 @@ class NvmlCudaPlatform(CudaPlatformBase):
     @classmethod
     @with_nvml_context
     def log_warnings(cls):
+        cls._warn_if_device_arch_not_compiled()
         device_ids: int = pynvml.nvmlDeviceGetCount()
         if device_ids > 1:
             device_names = [cls._get_physical_device_name(i) for i in range(device_ids)]