A high-throughput and memory-efficient inference and serving engine for LLMs
[Build] Fix CUDA arch build coverage gaps (#45277)
Signed-off-by: Shengqi Chen <harry-chen@outlook.com> Co-authored-by: Xin Li <xinli-sw@users.noreply.github.com> Co-authored-by: ShawRong <ShawRong@users.noreply.github.com> Co-authored-by: Change72 <Change72@users.noreply.github.com>
Shengqi Chen committed
Jun 14, 2026 at 13:09 UTC
78e7293bb157498d23780891cc4ef365a31772b6
14 files changed
+239
-131
.buildkite/release-pipeline.yaml
+16
-3
@@ -1,12 +1,25 @@
1
# CUDA architecture lists — following PyTorch RELEASE.md
2
# (https://github.com/pytorch/pytorch/blob/main/RELEASE.md)
3
# SM86 included for broader Ampere coverage; SM89 for marlin fp8 support
4
+# These requested arches are filtered by CMake's CUDA_SUPPORTED_ARCHS before
5
+# per-kernel arch selection. Do not add +PTX here: top-level +PTX is stripped
6
+# during that filtering, so kernels that need PTX must request it locally.
7
env:
5
- CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
6
- # aarch64 only architectures: 8.7 for Orin, 11.0 for Thor (since CUDA 13)
7
- CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0+PTX"
8
+ # for CUDA >=13, sm_100+ targets have family specifiers (see CMakeLists.txt)
9
+ # so targets like 10.3 and 12.1 are automatically supported with this list
10
+ CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
11
+ # aarch64-only targets: Orin (8.7), Thor (11.0, CUDA 13+)
12
+ CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0"
13
+
14
+ # for CUDA <13, we need to specify all needed targets
15
+ # some targets (10.3, 12.1) are skipped to limit the wheel size (< 500MB)
16
+ # please use CUDA 13 wheels or compile yourself on these new devices
17
CUDA_ARCH_X86_CU129: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
18
CUDA_ARCH_AARCH64_CU129: "8.0 8.7 8.9 9.0 10.0 12.0"
19
+
20
+ # pre-built mooncake wheels
21
+ # the manylinux_2_35 wheel has compatibility issue on Ubuntu 24.04
22
+ # so we use different wheels for the time being
23
MOONCAKE_WHEEL_AARCH64_2_35: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_aarch64.whl"
24
MOONCAKE_WHEEL_AARCH64_2_39: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_39_aarch64.whl"
25
MOONCAKE_WHEEL_X86_64: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_x86_64.whl"
.github/workflows/scripts/build.sh
+5
-2
@@ -9,7 +9,7 @@ PATH=${cuda_home}/bin:$PATH
9
LD_LIBRARY_PATH=${cuda_home}/lib64:$LD_LIBRARY_PATH
10
11
# Install requirements
12
-if [ "$(echo $2 | cut -d. -f1)" = "12" ]; then
12
+if [ "$(echo "$2" | cut -d. -f1)" = "12" ]; then
13
sed -i 's/^nvidia-cutlass-dsl\[cu13\]>=/nvidia-cutlass-dsl>=/' requirements/cuda.txt
14
fi
15
$python_executable -m pip install -r requirements/build/cuda.txt -r requirements/cuda.txt
@@ -17,7 +17,10 @@ $python_executable -m pip install -r requirements/build/cuda.txt -r requirements
17
# Limit the number of parallel jobs to avoid OOM
18
export MAX_JOBS=1
19
# Make sure release wheels are built for the following architectures
20
-export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
20
+# Do not add +PTX here: vLLM filters torch's top-level PTX flag when it
21
+# converts global gencode flags into per-kernel arch lists. If a specific
22
+# kernel needs PTX, add +PTX to that kernel's CMake arch list instead.
23
+export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0"
24
25
bash tools/check_repo.sh
26
CMakeLists.txt
+95
-88
@@ -190,6 +190,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
190
# the set of architectures we want to compile for and remove the from the
191
# CMAKE_CUDA_FLAGS so that they are not applied globally.
192
#
193
+ # `+PTX` in TORCH_CUDA_ARCH_LIST is not preserved here. It is emitted by torch
194
+ # as `code=compute_*`, while extract_unique_cuda_archs_ascending() records only
195
+ # `arch=compute_*`. If a kernel really needs PTX, add `+PTX` to that kernel's
196
+ # component-specific arch list below.
197
+ #
198
clear_cuda_arches(CUDA_ARCH_FLAGS)
199
extract_unique_cuda_archs_ascending(CUDA_ARCHS "${CUDA_ARCH_FLAGS}")
200
message(STATUS "CUDA target architectures: ${CUDA_ARCHS}")
@@ -365,13 +370,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
370
cuda_archs_loose_intersection(ES_MXFP8_GROUPED_MM_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
371
endif()
372
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND ES_MXFP8_GROUPED_MM_ARCHS)
368
- set(SRCS
373
+ set(ES_MXFP8_GROUPED_MM_SRCS
374
"csrc/libtorch_stable/moe/mxfp8_moe/cutlass_mxfp8_grouped_mm.cu"
375
"csrc/libtorch_stable/moe/mxfp8_moe/mxfp8_experts_quant.cu")
376
set_gencode_flags_for_srcs(
372
- SRCS "${SRCS}"
377
+ SRCS "${ES_MXFP8_GROUPED_MM_SRCS}"
378
CUDA_ARCHS "${ES_MXFP8_GROUPED_MM_ARCHS}")
374
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
379
+ list(APPEND VLLM_STABLE_EXT_SRC "${ES_MXFP8_GROUPED_MM_SRCS}")
380
list(APPEND VLLM_GPU_FLAGS "-DENABLE_ES_MXFP8_GROUPED_MM_SM100=1")
381
message(STATUS "Building ES MXFP8 grouped kernels for archs: ${ES_MXFP8_GROUPED_MM_ARCHS}")
382
else()
@@ -676,16 +681,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
681
682
# DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
683
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
679
- cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
684
+ cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
685
else()
681
- cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
686
+ cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
687
endif()
688
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
684
- set(SRCS "csrc/libtorch_stable/dsv3_fused_a_gemm.cu")
689
+ set(DSV3_FUSED_A_GEMM_SRCS "csrc/libtorch_stable/dsv3_fused_a_gemm.cu")
690
set_gencode_flags_for_srcs(
686
- SRCS "${SRCS}"
691
+ SRCS "${DSV3_FUSED_A_GEMM_SRCS}"
692
CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
688
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
693
+ list(APPEND VLLM_STABLE_EXT_SRC "${DSV3_FUSED_A_GEMM_SRCS}")
694
message(STATUS "Building dsv3_fused_a_gemm for archs: ${DSV3_FUSED_A_GEMM_ARCHS}")
695
else()
696
message(STATUS "Not building dsv3_fused_a_gemm as no compatible archs found "
@@ -695,13 +700,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
700
# FP32 router GEMM (H=3072, E=256, M<=32). Requires SM90+ and CUDA >= 12.0.
701
cuda_archs_sm90plus(FP32_ROUTER_GEMM_ARCHS "${CUDA_ARCHS}")
702
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND FP32_ROUTER_GEMM_ARCHS)
698
- set(SRCS
703
+ set(FP32_ROUTER_GEMM_SRCS
704
"csrc/libtorch_stable/fp32_router_gemm_entry.cu"
705
"csrc/libtorch_stable/fp32_router_gemm.cu")
706
set_gencode_flags_for_srcs(
702
- SRCS "${SRCS}"
707
+ SRCS "${FP32_ROUTER_GEMM_SRCS}"
708
CUDA_ARCHS "${FP32_ROUTER_GEMM_ARCHS}")
704
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
709
+ list(APPEND VLLM_STABLE_EXT_SRC "${FP32_ROUTER_GEMM_SRCS}")
710
message(STATUS "Building fp32_router_gemm for archs: ${FP32_ROUTER_GEMM_ARCHS}")
711
else()
712
message(STATUS "Not building fp32_router_gemm as no compatible archs found "
@@ -711,13 +716,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
716
# Only build AllSpark kernels if we are building for at least some compatible archs.
717
cuda_archs_loose_intersection(ALLSPARK_ARCHS "8.0;8.6;8.7;8.9" "${CUDA_ARCHS}")
718
if (ALLSPARK_ARCHS)
714
- set(SRCS
719
+ set(ALLSPARK_SRCS
720
"csrc/libtorch_stable/quantization/gptq_allspark/allspark_repack.cu"
721
"csrc/libtorch_stable/quantization/gptq_allspark/allspark_qgemm_w8a16.cu")
722
set_gencode_flags_for_srcs(
718
- SRCS "${SRCS}"
723
+ SRCS "${ALLSPARK_SRCS}"
724
CUDA_ARCHS "${ALLSPARK_ARCHS}")
720
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
725
+ list(APPEND VLLM_STABLE_EXT_SRC "${ALLSPARK_SRCS}")
726
message(STATUS "Building AllSpark kernels for archs: ${ALLSPARK_ARCHS}")
727
else()
728
message(STATUS "Not building AllSpark kernels as no compatible archs found"
@@ -732,16 +737,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
737
# CUDA 12.0 or later
738
cuda_archs_loose_intersection(SCALED_MM_ARCHS "9.0a;" "${CUDA_ARCHS}")
739
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND SCALED_MM_ARCHS)
735
- set(SRCS
740
+ set(SCALED_MM_SM90_SRCS
741
"csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c3x_sm90.cu"
742
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_fp8.cu"
743
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm90_int8.cu"
744
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_azp_sm90_int8.cu"
745
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm90_fp8.cu")
746
set_gencode_flags_for_srcs(
742
- SRCS "${SRCS}"
747
+ SRCS "${SCALED_MM_SM90_SRCS}"
748
CUDA_ARCHS "${SCALED_MM_ARCHS}")
744
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
749
+ list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_SM90_SRCS}")
750
list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_SM90=1")
751
# Let scaled_mm_c2x know it doesn't need to build these arches
752
list(APPEND SCALED_MM_3X_ARCHS "${SCALED_MM_ARCHS}")
@@ -767,15 +772,15 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
772
cuda_archs_loose_intersection(SCALED_MM_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
773
endif()
774
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
770
- set(SRCS
775
+ set(SCALED_MM_SM120_SRCS
776
"csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c3x_sm120.cu"
777
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8.cu"
778
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm120_fp8.cu"
779
)
780
set_gencode_flags_for_srcs(
776
- SRCS "${SRCS}"
781
+ SRCS "${SCALED_MM_SM120_SRCS}"
782
CUDA_ARCHS "${SCALED_MM_ARCHS}")
778
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
783
+ list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_SM120_SRCS}")
784
list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_SM120=1")
785
# Let scaled_mm_c2x know it doesn't need to build these arches
786
list(APPEND SCALED_MM_3X_ARCHS "${SCALED_MM_ARCHS}")
@@ -801,15 +806,15 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
806
cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
807
endif()
808
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
804
- set(SRCS
809
+ set(SCALED_MM_SM100_SRCS
810
"csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c3x_sm100.cu"
811
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm100_fp8.cu"
812
"csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm100_fp8.cu"
813
)
814
set_gencode_flags_for_srcs(
810
- SRCS "${SRCS}"
815
+ SRCS "${SCALED_MM_SM100_SRCS}"
816
CUDA_ARCHS "${SCALED_MM_ARCHS}")
812
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
817
+ list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_SM100_SRCS}")
818
list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_SM100=1")
819
# Let scaled_mm_c2x know it doesn't need to build these arches
820
list(APPEND SCALED_MM_3X_ARCHS "${SCALED_MM_ARCHS}")
@@ -835,11 +840,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
840
# subtract out the archs that are already built for 3x
841
list(REMOVE_ITEM SCALED_MM_2X_ARCHS ${SCALED_MM_3X_ARCHS})
842
if (SCALED_MM_2X_ARCHS)
838
- set(SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c2x.cu")
843
+ set(SCALED_MM_C2X_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_c2x.cu")
844
set_gencode_flags_for_srcs(
840
- SRCS "${SRCS}"
845
+ SRCS "${SCALED_MM_C2X_SRCS}"
846
CUDA_ARCHS "${SCALED_MM_2X_ARCHS}")
842
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
847
+ list(APPEND VLLM_STABLE_EXT_SRC "${SCALED_MM_C2X_SRCS}")
848
list(APPEND VLLM_GPU_FLAGS "-DENABLE_SCALED_MM_C2X=1")
849
message(STATUS "Building scaled_mm_c2x for archs: ${SCALED_MM_2X_ARCHS}")
850
else()
@@ -861,11 +866,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
866
# if it's possible to compile MoE kernels that use its output.
867
cuda_archs_loose_intersection(SCALED_MM_ARCHS "9.0a" "${CUDA_ARCHS}")
868
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.3 AND SCALED_MM_ARCHS)
864
- set(SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm90.cu")
869
+ set(CUTLASS_MOE_SM90_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm90.cu")
870
set_gencode_flags_for_srcs(
866
- SRCS "${SRCS}"
871
+ SRCS "${CUTLASS_MOE_SM90_SRCS}"
872
CUDA_ARCHS "${SCALED_MM_ARCHS}")
868
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
873
+ list(APPEND VLLM_STABLE_EXT_SRC "${CUTLASS_MOE_SM90_SRCS}")
874
list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM90=1")
875
message(STATUS "Building grouped_mm_c3x for archs: ${SCALED_MM_ARCHS}")
876
else()
@@ -880,16 +885,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
885
endif()
886
887
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
883
- cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
888
+ cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0f" "${CUDA_ARCHS}")
889
else()
885
- cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
890
+ cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0a;10.3a" "${CUDA_ARCHS}")
891
endif()
892
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
888
- set(SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm100.cu")
893
+ set(CUTLASS_MOE_SM100_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/grouped_mm_c3x_sm100.cu")
894
set_gencode_flags_for_srcs(
890
- SRCS "${SRCS}"
895
+ SRCS "${CUTLASS_MOE_SM100_SRCS}"
896
CUDA_ARCHS "${SCALED_MM_ARCHS}")
892
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
897
+ list(APPEND VLLM_STABLE_EXT_SRC "${CUTLASS_MOE_SM100_SRCS}")
898
list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM100=1")
899
message(STATUS "Building grouped_mm_c3x for archs: ${SCALED_MM_ARCHS}")
900
else()
@@ -910,11 +915,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
915
cuda_archs_loose_intersection(CUTLASS_MOE_DATA_ARCHS "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
916
endif()
917
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.3 AND CUTLASS_MOE_DATA_ARCHS)
913
- set(SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/moe_data.cu")
918
+ set(CUTLASS_MOE_DATA_SRCS "csrc/libtorch_stable/quantization/w8a8/cutlass/moe/moe_data.cu")
919
set_gencode_flags_for_srcs(
915
- SRCS "${SRCS}"
920
+ SRCS "${CUTLASS_MOE_DATA_SRCS}"
921
CUDA_ARCHS "${CUTLASS_MOE_DATA_ARCHS}")
917
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
922
+ list(APPEND VLLM_STABLE_EXT_SRC "${CUTLASS_MOE_DATA_SRCS}")
923
message(STATUS "Building moe_data for archs: ${CUTLASS_MOE_DATA_ARCHS}")
924
else()
925
if (NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.3 AND CUTLASS_MOE_DATA_ARCHS)
@@ -931,71 +936,66 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
936
# FP4/NVFP4 kernels (moved from _C to _C_stable_libtorch)
937
#
938
934
- # The nvfp4_scaled_mm_sm120 kernels for Blackwell SM12x require
935
- # CUDA 12.8 or later
939
+ # SM12x FP4 kernels. These share some generic NVFP4 quantization entry
940
+ # sources with the SM10x/11x block below; set_gencode_flags_for_srcs appends
941
+ # per-source flags, so shared files accumulate both SM12x and SM10x/11x
942
+ # gencodes when both families are requested.
943
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
937
- cuda_archs_loose_intersection(FP4_ARCHS "12.0f" "${CUDA_ARCHS}")
944
+ cuda_archs_loose_intersection(FP4_SM120_ARCHS "12.0f" "${CUDA_ARCHS}")
945
else()
939
- cuda_archs_loose_intersection(FP4_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
946
+ cuda_archs_loose_intersection(FP4_SM120_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
947
endif()
941
- if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_ARCHS)
942
- set(SRCS
948
+ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_SM120_ARCHS)
949
+ set(FP4_SM120_SRCS
950
"csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu"
951
"csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
952
"csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
953
"csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_kernels.cu"
947
- "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu")
948
- set_gencode_flags_for_srcs(
949
- SRCS "${SRCS}"
950
- CUDA_ARCHS "${FP4_ARCHS}")
951
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
952
- set(NVFP4_KV_SRC "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
954
+ "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
955
+ "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
956
set_gencode_flags_for_srcs(
954
- SRCS "${NVFP4_KV_SRC}"
955
- CUDA_ARCHS "${FP4_ARCHS}")
956
- list(APPEND VLLM_STABLE_EXT_SRC "${NVFP4_KV_SRC}")
957
+ SRCS "${FP4_SM120_SRCS}"
958
+ CUDA_ARCHS "${FP4_SM120_ARCHS}")
959
+ list(APPEND VLLM_STABLE_EXT_SRC "${FP4_SM120_SRCS}")
960
target_compile_definitions(_C PRIVATE ENABLE_NVFP4_SM120=1)
961
list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM120=1")
962
list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM120=1")
960
- message(STATUS "Building NVFP4 for archs: ${FP4_ARCHS}")
963
+ message(STATUS "Building SM12x NVFP4 for archs: ${FP4_SM120_ARCHS}")
964
else()
962
- message(STATUS "Not building NVFP4 as no compatible archs were found.")
963
- # clear FP4_ARCHS
964
- set(FP4_ARCHS)
965
+ message(STATUS "Not building SM12x NVFP4 as no compatible archs were found.")
966
endif()
967
967
- # FP4 Archs and flags
968
+ # SM10x/11x FP4 kernels. MXFP4 experts quantization is currently compiled
969
+ # only in this block; SM12x has separate NVFP4 matmul/MoE kernels above.
970
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
969
- cuda_archs_loose_intersection(FP4_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
971
+ cuda_archs_loose_intersection(FP4_SM100_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
972
else()
971
- cuda_archs_loose_intersection(FP4_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
973
+ cuda_archs_loose_intersection(FP4_SM100_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
974
endif()
973
- if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_ARCHS)
974
- set(SRCS
975
+ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_SM100_ARCHS)
976
+ set(FP4_SM100_SRCS
977
"csrc/libtorch_stable/quantization/fp4/nvfp4_quant_kernels.cu"
978
"csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
979
"csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
980
"csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
981
"csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
982
"csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu"
981
- "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu")
982
- set_gencode_flags_for_srcs(
983
- SRCS "${SRCS}"
984
- CUDA_ARCHS "${FP4_ARCHS}")
985
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
986
- set(NVFP4_KV_SRC "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
983
+ "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu"
984
+ "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu")
985
+ if(NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9)
986
+ message(STATUS
987
+ "Building mxfp4_experts_quant unsupported stubs because CUDA compiler version is not >= 12.9 (found ${CMAKE_CUDA_COMPILER_VERSION}).")
988
+ endif()
989
set_gencode_flags_for_srcs(
988
- SRCS "${NVFP4_KV_SRC}"
989
- CUDA_ARCHS "${FP4_ARCHS}")
990
- list(APPEND VLLM_STABLE_EXT_SRC "${NVFP4_KV_SRC}")
990
+ SRCS "${FP4_SM100_SRCS}"
991
+ CUDA_ARCHS "${FP4_SM100_ARCHS}")
992
+ list(APPEND VLLM_STABLE_EXT_SRC "${FP4_SM100_SRCS}")
993
target_compile_definitions(_C PRIVATE ENABLE_NVFP4_SM100=1)
994
list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM100=1")
995
list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM100=1")
994
- message(STATUS "Building NVFP4 for archs: ${FP4_ARCHS}")
996
+ message(STATUS "Building SM10x/11x NVFP4/MXFP4 for archs: ${FP4_SM100_ARCHS}")
997
else()
996
- message(STATUS "Not building NVFP4 as no compatible archs were found.")
997
- # clear FP4_ARCHS
998
- set(FP4_ARCHS)
998
+ message(STATUS "Not building SM10x/11x NVFP4/MXFP4 as no compatible archs were found.")
999
endif()
1000
1001
#
@@ -1005,17 +1005,17 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
1005
# Only build W4A8 kernels if we are building for something compatible with sm90a
1006
cuda_archs_loose_intersection(W4A8_ARCHS "9.0a" "${CUDA_ARCHS}")
1007
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND W4A8_ARCHS)
1008
- set(SRCS
1008
+ set(W4A8_SRCS
1009
"csrc/libtorch_stable/quantization/cutlass_w4a8/w4a8_mm_entry.cu"
1010
"csrc/libtorch_stable/quantization/cutlass_w4a8/w4a8_grouped_mm_entry.cu"
1011
"csrc/libtorch_stable/quantization/cutlass_w4a8/w4a8_utils.cu"
1012
)
1013
1014
set_gencode_flags_for_srcs(
1015
- SRCS "${SRCS}"
1015
+ SRCS "${W4A8_SRCS}"
1016
CUDA_ARCHS "${W4A8_ARCHS}")
1017
1018
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
1018
+ list(APPEND VLLM_STABLE_EXT_SRC "${W4A8_SRCS}")
1019
1020
message(STATUS "Building W4A8 kernels for archs: ${W4A8_ARCHS}")
1021
else()
@@ -1031,22 +1031,24 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
1031
endif()
1032
endif()
1033
1034
- # CUTLASS MLA Archs and flags
1034
+ # CUTLASS MLA Archs and flags.
1035
+ # Runtime dispatch is gated in
1036
+ # vllm/v1/attention/backends/mla/cutlass_mla.py.
1037
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
1036
- cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
1038
+ cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
1039
else()
1038
- cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
1040
+ cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
1041
endif()
1042
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
1041
- set(SRCS
1043
+ set(CUTLASS_MLA_SRCS
1044
"csrc/libtorch_stable/attention/mla/sm100_cutlass_mla_kernel.cu")
1045
set_gencode_flags_for_srcs(
1044
- SRCS "${SRCS}"
1046
+ SRCS "${CUTLASS_MLA_SRCS}"
1047
CUDA_ARCHS "${MLA_ARCHS}")
1046
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
1048
+ list(APPEND VLLM_STABLE_EXT_SRC "${CUTLASS_MLA_SRCS}")
1049
list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
1050
# Add MLA-specific include directories only to MLA source files
1049
- set_source_files_properties(${SRCS}
1051
+ set_source_files_properties(${CUTLASS_MLA_SRCS}
1052
PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
1053
message(STATUS "Building CUTLASS MLA for archs: ${MLA_ARCHS}")
1054
else()
@@ -1058,11 +1060,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
1060
# Hadacore kernels
1061
cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
1062
if(HADACORE_ARCHS)
1061
- set(SRCS "csrc/libtorch_stable/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
1063
+ set(HADACORE_SRCS "csrc/libtorch_stable/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
1064
set_gencode_flags_for_srcs(
1063
- SRCS "${SRCS}"
1065
+ SRCS "${HADACORE_SRCS}"
1066
CUDA_ARCHS "${HADACORE_ARCHS}")
1065
- list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
1067
+ list(APPEND VLLM_STABLE_EXT_SRC "${HADACORE_SRCS}")
1068
message(STATUS "Building hadacore")
1069
endif()
1070
@@ -1070,6 +1072,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
1072
endif()
1073
1074
message(STATUS "Enabling C_stable extension.")
1075
+ list(REMOVE_DUPLICATES VLLM_STABLE_EXT_SRC)
1076
define_extension_target(
1077
_C_stable_libtorch
1078
DESTINATION vllm
@@ -1174,7 +1177,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
1177
# - sm80 doesn't support fp8 computation
1178
# - sm90 and sm100 don't support QMMA.16832.F32.E4M3.E4M3 SAAS instruction
1179
# so we only enable fp8 computation for SM89 (e.g. RTX 40x0) and 12.0 (e.g. RTX 50x0)
1177
- cuda_archs_loose_intersection(MARLIN_MOE_FP8_ARCHS "8.9;12.0;12.1" "${CUDA_ARCHS}")
1180
+ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
1181
+ cuda_archs_loose_intersection(MARLIN_MOE_FP8_ARCHS "8.9;12.0f" "${CUDA_ARCHS}")
1182
+ else()
1183
+ cuda_archs_loose_intersection(MARLIN_MOE_FP8_ARCHS "8.9;12.0a;12.1a" "${CUDA_ARCHS}")
1184
+ endif()
1185
# moe marlin arches for other files
1186
cuda_archs_loose_intersection(MARLIN_MOE_OTHER_ARCHS "7.5;8.0+PTX" "${CUDA_ARCHS}")
1187
if (MARLIN_MOE_OTHER_ARCHS)
cmake/external_projects/qutlass.cmake
+22
-10
@@ -32,21 +32,33 @@ endif()
32
message(STATUS "[QUTLASS] QuTLASS is available at ${qutlass_SOURCE_DIR}")
33
34
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
35
- cuda_archs_loose_intersection(QUTLASS_ARCHS "10.0f;12.0f" "${CUDA_ARCHS}")
35
+ cuda_archs_loose_intersection(QUTLASS_SM120_ARCHS "12.0f" "${CUDA_ARCHS}")
36
+ cuda_archs_loose_intersection(QUTLASS_SM100_ARCHS "10.0f" "${CUDA_ARCHS}")
37
else()
37
- cuda_archs_loose_intersection(QUTLASS_ARCHS "12.0a;12.1a;10.0a;10.3a" "${CUDA_ARCHS}")
38
+ cuda_archs_loose_intersection(QUTLASS_SM120_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
39
+ cuda_archs_loose_intersection(QUTLASS_SM100_ARCHS "10.0a;10.3a" "${CUDA_ARCHS}")
40
endif()
41
40
-if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
41
-
42
- if(QUTLASS_ARCHS MATCHES "10\\.(0a|3a|0f)")
43
- set(QUTLASS_TARGET_CC 100)
44
- elseif(QUTLASS_ARCHS MATCHES "12\\.[01][af]?")
45
- set(QUTLASS_TARGET_CC 120)
46
- else()
47
- message(FATAL_ERROR "[QUTLASS] internal error parsing CUDA_ARCHS='${QUTLASS_ARCHS}'.")
42
+# QUTLASS uses TARGET_CUDA_ARCH as a single preprocessor selector for all its
43
+# sources. Do not compile a mixed SM100/SM120 arch list with one selector; prefer
44
+# SM100 when both families are requested because that is the primary deployed
45
+# target for this extension today.
46
+if(QUTLASS_SM100_ARCHS)
47
+ set(QUTLASS_ARCHS "${QUTLASS_SM100_ARCHS}")
48
+ set(QUTLASS_TARGET_CC 100)
49
+ if(QUTLASS_SM120_ARCHS)
50
+ message(WARNING
51
+ "[QUTLASS] Both SM100 and SM120 archs were requested; selecting SM100 "
52
+ "because TARGET_CUDA_ARCH is a single compile-time selector.")
53
endif()
54
+elseif(QUTLASS_SM120_ARCHS)
55
+ set(QUTLASS_ARCHS "${QUTLASS_SM120_ARCHS}")
56
+ set(QUTLASS_TARGET_CC 120)
57
+else()
58
+ set(QUTLASS_ARCHS)
59
+endif()
60
61
+if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
62
set(QUTLASS_SOURCES
63
${qutlass_SOURCE_DIR}/qutlass/csrc/bindings.cpp
64
${qutlass_SOURCE_DIR}/qutlass/csrc/gemm.cu
cmake/utils.cmake
+2
-2
@@ -487,9 +487,9 @@ endfunction()
487
488
function(cuda_archs_sm90plus OUT_CUDA_ARCHS TGT_CUDA_ARCHS)
489
if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
490
- cuda_archs_loose_intersection(_archs "9.0a;10.0f;11.0f" "${TGT_CUDA_ARCHS}")
490
+ cuda_archs_loose_intersection(_archs "9.0a;10.0f;11.0f;12.0f" "${TGT_CUDA_ARCHS}")
491
else()
492
- cuda_archs_loose_intersection(_archs "9.0a;10.0a;10.1a;10.3a" "${TGT_CUDA_ARCHS}")
492
+ cuda_archs_loose_intersection(_archs "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${TGT_CUDA_ARCHS}")
493
endif()
494
set(${OUT_CUDA_ARCHS} ${_archs} PARENT_SCOPE)
495
endfunction()
csrc/libtorch_stable/cuda_vec_utils.cuh
+1
-1
@@ -21,7 +21,7 @@
21
// together enable 256-bit (v8.u32) PTX load/store instructions.
22
// Use for PTX instruction selection with architecture fallback paths.
23
#if !defined(USE_ROCM) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1000 && \
24
- defined(CUDA_VERSION) && CUDA_VERSION >= 12090
24
+ defined(CUDART_VERSION) && CUDART_VERSION >= 12090
25
#define VLLM_256B_PTX_ENABLED 1
26
#else
27
#define VLLM_256B_PTX_ENABLED 0
csrc/libtorch_stable/moe/dsv3_router_gemm_entry.cu
+1
-2
@@ -144,8 +144,7 @@ void dsv3_router_gemm(
144
"output must be float32 or bf16");
145
146
const int sm = getSMVersion();
147
- STD_TORCH_CHECK(sm >= 90 && sm <= 103,
148
- "required SM_103 >= CUDA ARCH >= SM_90");
147
+ STD_TORCH_CHECK(sm >= 90, "required CUDA ARCH >= SM_90");
148
149
const cudaStream_t stream = get_current_cuda_stream(mat_a.get_device_index());
150
csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu
+63
-12
@@ -27,15 +27,24 @@
27
#include <torch/csrc/stable/tensor.h>
28
#include "libtorch_stable/torch_utils.h"
29
#include "libtorch_stable/dispatch_utils.h"
30
+#include "libtorch_stable/cutlass_extensions/common.hpp"
31
#include "../../cuda_vec_utils.cuh"
32
#include "cuda_utils.h"
33
34
#include "nvfp4_utils.cuh"
35
+
36
+#if defined(CUDART_VERSION) && CUDART_VERSION >= 12090
37
+ #define VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED 1
38
static_assert(CVT_FP4_ELTS_PER_THREAD == 16,
39
"MXFP4 experts quant requires PACK16 mode (CUDA >= 12.9)");
40
+#else
41
+ #define VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED 0
42
+#endif
43
44
#include "libtorch_stable/launch_bounds_utils.h"
45
46
+#if VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED
47
+
48
namespace vllm {
49
50
// MXFP4 block size constants
@@ -104,7 +113,7 @@ __global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
113
&input_offset_by_experts[chunk_start + 12]));
114
local_offsets[16] = __ldca(&input_offset_by_experts[chunk_start + 16]);
115
107
-#pragma unroll
116
+ #pragma unroll
117
for (int i = 0; i < 16; i++) {
118
if (rowIdx >= local_offsets[i] && rowIdx < local_offsets[i + 1]) {
119
rowIdx_in_expert = rowIdx - local_offsets[i];
@@ -309,14 +318,14 @@ void mxfp4_quant_impl(void* output, void* output_scale, void* input,
318
319
} // namespace vllm
320
312
-/*Quantization entry for mxfp4 experts quantization*/
313
-#define CHECK_TH_CUDA(x, m) \
314
- STD_TORCH_CHECK(x.is_cuda(), m, "must be a CUDA tensor")
315
-#define CHECK_CONTIGUOUS(x, m) \
316
- STD_TORCH_CHECK(x.is_contiguous(), m, "must be contiguous")
317
-#define CHECK_INPUT(x, m) \
318
- CHECK_TH_CUDA(x, m); \
319
- CHECK_CONTIGUOUS(x, m);
321
+ /*Quantization entry for mxfp4 experts quantization*/
322
+ #define CHECK_TH_CUDA(x, m) \
323
+ STD_TORCH_CHECK(x.is_cuda(), m, "must be a CUDA tensor")
324
+ #define CHECK_CONTIGUOUS(x, m) \
325
+ STD_TORCH_CHECK(x.is_contiguous(), m, "must be contiguous")
326
+ #define CHECK_INPUT(x, m) \
327
+ CHECK_TH_CUDA(x, m); \
328
+ CHECK_CONTIGUOUS(x, m);
329
330
constexpr auto HALF = torch::headeronly::ScalarType::Half;
331
constexpr auto BF16 = torch::headeronly::ScalarType::BFloat16;
@@ -364,12 +373,28 @@ static void validate_mxfp4_experts_quant_inputs(
373
STD_TORCH_CHECK(output_scale.size(1) * 4 == padded_k);
374
}
375
376
+#endif // VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED
377
+
378
+static bool mxfp4_experts_quant_sm_supported(int64_t cuda_device_capability) {
379
+#if VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED
380
+ return cuda_device_capability >= 100 && cuda_device_capability < 120;
381
+#else
382
+ return false;
383
+#endif
384
+}
385
+
386
void mxfp4_experts_quant(
387
torch::stable::Tensor& output, torch::stable::Tensor& output_scale,
388
torch::stable::Tensor const& input,
389
torch::stable::Tensor const& input_offset_by_experts,
390
torch::stable::Tensor const& output_scale_offset_by_experts,
391
int64_t n_experts) {
392
+#if VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED
393
+ int32_t sm = get_sm_version_num();
394
+ STD_TORCH_CHECK(mxfp4_experts_quant_sm_supported(sm),
395
+ "No compiled MXFP4 experts quant kernel for SM ", sm,
396
+ ". Recompile with SM10x/11x FP4 support and CUDA >= 12.9.");
397
+
398
auto m_topk = input.size(0);
399
auto k = input.size(1);
400
@@ -390,6 +415,10 @@ void mxfp4_experts_quant(
415
output_scale_offset_by_experts.data_ptr(), m_topk, k, n_experts,
416
stream);
417
});
418
+#else
419
+ STD_TORCH_CHECK_NOT_IMPLEMENTED(false,
420
+ "MXFP4 experts quant requires CUDA >= 12.9.");
421
+#endif
422
}
423
424
void silu_and_mul_mxfp4_experts_quant(
@@ -398,6 +427,12 @@ void silu_and_mul_mxfp4_experts_quant(
427
torch::stable::Tensor const& input_offset_by_experts,
428
torch::stable::Tensor const& output_scale_offset_by_experts,
429
int64_t n_experts) {
430
+#if VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED
431
+ int32_t sm = get_sm_version_num();
432
+ STD_TORCH_CHECK(mxfp4_experts_quant_sm_supported(sm),
433
+ "No compiled SiLU+Mul MXFP4 experts quant kernel for SM ", sm,
434
+ ". Recompile with SM10x/11x FP4 support and CUDA >= 12.9.");
435
+
436
auto m_topk = input.size(0);
437
auto k_times_2 = input.size(1);
438
STD_TORCH_CHECK(k_times_2 % 2 == 0, "input width must be even (gate || up)");
@@ -420,13 +455,29 @@ void silu_and_mul_mxfp4_experts_quant(
455
output_scale_offset_by_experts.data_ptr(), m_topk, k, n_experts,
456
stream);
457
});
458
+#else
459
+ STD_TORCH_CHECK_NOT_IMPLEMENTED(
460
+ false, "SiLU+Mul MXFP4 experts quant requires CUDA >= 12.9.");
461
+#endif
462
}
463
425
-// Registered here (not torch_bindings.cpp) because VLLM_GPU_FLAGS is applied
426
-// only under COMPILE_LANGUAGE:CUDA, so ENABLE_NVFP4_SM100 is invisible to
427
-// .cpp files and cannot gate the registration from there.
464
+bool mxfp4_experts_quant_supported(int64_t cuda_device_capability) {
465
+ return mxfp4_experts_quant_sm_supported(cuda_device_capability);
466
+}
467
+
468
+STABLE_TORCH_LIBRARY_FRAGMENT(_C, m) {
469
+ m.def("mxfp4_experts_quant_supported(int cuda_device_capability) -> bool");
470
+}
471
+
472
+// Registered here so the CUDA 12.8 stub and CUDA 12.9+ implementation stay
473
+// tied to the same translation unit.
474
STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
475
m.impl("mxfp4_experts_quant", TORCH_BOX(&mxfp4_experts_quant));
476
m.impl("silu_and_mul_mxfp4_experts_quant",
477
TORCH_BOX(&silu_and_mul_mxfp4_experts_quant));
478
}
479
+
480
+STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, m) {
481
+ m.impl("mxfp4_experts_quant_supported",
482
+ TORCH_BOX(&mxfp4_experts_quant_supported));
483
+}
csrc/libtorch_stable/quantization/fp4/nvfp4_utils.cuh
+4
-4
@@ -22,15 +22,15 @@
22
23
#include "../../cuda_vec_utils.cuh"
24
25
-#if defined(NVFP4_ENABLE_ELTS16) && defined(CUDA_VERSION) && \
26
- CUDA_VERSION >= 12090
25
+#if defined(NVFP4_ENABLE_ELTS16) && defined(CUDART_VERSION) && \
26
+ CUDART_VERSION >= 12090
27
#define ELTS_PER_THREAD 16
28
+ #define CVT_FP4_PACK16 1
29
constexpr int CVT_FP4_ELTS_PER_THREAD = 16;
29
-constexpr bool CVT_FP4_PACK16 = true;
30
#else
31
#define ELTS_PER_THREAD 8
32
+ #define CVT_FP4_PACK16 0
33
constexpr int CVT_FP4_ELTS_PER_THREAD = 8;
33
-constexpr bool CVT_FP4_PACK16 = false;
34
#endif
35
36
constexpr int CVT_FP4_SF_VEC_SIZE = 16;
csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu
+9
-3
@@ -1,3 +1,4 @@
1
+#include <cuda.h>
2
#include <cudaTypedefs.h>
3
4
#include <torch/csrc/stable/tensor.h>
@@ -174,15 +175,20 @@ bool cutlass_scaled_mm_supports_block_fp8(int64_t cuda_device_capability) {
175
176
bool cutlass_group_gemm_supported(int64_t cuda_device_capability) {
177
// CUTLASS grouped FP8 kernels need at least CUDA 12.3 and SM90 (Hopper)
177
- // or CUDA 12.8 and SM100 (Blackwell)
178
+ // or CUDA 12.8 and SM100 (Blackwell). Only report archs that have an
179
+ // actual cutlass_moe_mm dispatch compiled into this file.
180
181
#if defined CUDA_VERSION
180
- if (cuda_device_capability >= 100) {
182
+ #if defined ENABLE_CUTLASS_MOE_SM100 && ENABLE_CUTLASS_MOE_SM100
183
+ if (cuda_device_capability >= 100 && cuda_device_capability < 110) {
184
return CUDA_VERSION >= 12080;
185
}
183
- if (cuda_device_capability >= 90) {
186
+ #endif
187
+ #if defined ENABLE_CUTLASS_MOE_SM90 && ENABLE_CUTLASS_MOE_SM90
188
+ if (cuda_device_capability >= 90 && cuda_device_capability < 100) {
189
return CUDA_VERSION >= 12030;
190
}
191
+ #endif
192
#endif
193
194
return false;
docker/Dockerfile
+6
-2
@@ -261,7 +261,10 @@ RUN --mount=type=cache,target=/opt/uv/cache \
261
# Explicitly set the list to avoid issues with torch 2.2
262
# See https://github.com/pytorch/pytorch/pull/123243
263
# From versions.json: .torch.cuda_arch_list
264
-ARG torch_cuda_arch_list='7.5 8.0 8.6 8.9 9.0 10.0 11.0 12.0+PTX'
264
+# Do not add +PTX here: vLLM filters torch's top-level PTX flag when it
265
+# converts global gencode flags into per-kernel arch lists. If a specific
266
+# kernel needs PTX, add +PTX to that kernel's CMake arch list instead.
267
+ARG torch_cuda_arch_list='7.5 8.0 8.6 8.9 9.0 10.0 11.0 12.0'
268
ENV TORCH_CUDA_ARCH_LIST=${torch_cuda_arch_list}
269
#################### BUILD BASE IMAGE ####################
270
@@ -1010,7 +1013,8 @@ ARG PIP_EXTRA_INDEX_URL UV_EXTRA_INDEX_URL
1013
ENV UV_HTTP_TIMEOUT=500
1014
1015
# install kv_connectors if requested
1013
-ARG torch_cuda_arch_list='7.5 8.0 8.6 8.9 9.0 10.0 11.0 12.0+PTX'
1016
+# Do not add +PTX here; see the main TORCH_CUDA_ARCH_LIST comment above.
1017
+ARG torch_cuda_arch_list='7.5 8.0 8.6 8.9 9.0 10.0 11.0 12.0'
1018
ENV TORCH_CUDA_ARCH_LIST=${torch_cuda_arch_list}
1019
RUN --mount=type=cache,target=/opt/uv/cache \
1020
--mount=type=bind,source=requirements/kv_connectors.txt,target=/tmp/kv_connectors.txt,ro \
docker/versions.json
+1
-1
@@ -35,7 +35,7 @@
35
"default": "false"
36
},
37
"TORCH_CUDA_ARCH_LIST": {
38
- "default": "7.5 8.0 8.6 8.9 9.0 10.0 11.0 12.0+PTX"
38
+ "default": "7.5 8.0 8.6 8.9 9.0 10.0 11.0 12.0"
39
},
40
"MAX_JOBS": {
41
"default": "2"
vllm/_custom_ops.py
+8
@@ -773,6 +773,14 @@ def cutlass_scaled_mm_supports_fp4(cuda_device_capability: int) -> bool:
773
return torch.ops._C.cutlass_scaled_mm_supports_fp4(cuda_device_capability)
774
775
776
+def mxfp4_experts_quant_supported(cuda_device_capability: int) -> bool:
777
+ try:
778
+ return torch.ops._C.mxfp4_experts_quant_supported(cuda_device_capability)
779
+ except AttributeError:
780
+ # Return False on builds where the CUDA helper is not available.
781
+ return False
782
+
783
+
784
def cutlass_scaled_fp4_mm(
785
a: torch.Tensor,
786
b: torch.Tensor,
vllm/model_executor/layers/fused_moe/experts/cutlass_moe.py
+6
-1
@@ -997,7 +997,12 @@ class CutlassExpertsMxfp4(mk.FusedMoEExpertsModular):
997
@staticmethod
998
def _supports_current_device() -> bool:
999
p = current_platform
1000
- return p.is_cuda() and p.is_device_capability_family(100)
1000
+ capability = p.get_device_capability()
1001
+ return (
1002
+ p.is_cuda()
1003
+ and capability is not None
1004
+ and ops.mxfp4_experts_quant_supported(capability.to_int())
1005
+ )
1006
1007
@staticmethod
1008
def _supports_no_act_and_mul() -> bool: