From 19932fdcdaa4b40fbdad93acd179054aba173140 Mon Sep 17 00:00:00 2001 From: ayrnb <641876696@qq.com> Date: Wed, 5 Aug 2026 09:11:12 +0000 Subject: [PATCH 1/2] [Kernel] Enable CUTLASS MXFP4 W4A4 MoE on SM12x The MXFP4 experts-quant and grouped-GEMM kernels were compiled and gated for SM10x/11x only, so MXFP4 W4A4 MoE fell back to Marlin W4A16 on GeForce / RTX PRO Blackwell. CUTLASS resolves an identical block-scaled scale-factor layout for arch::Sm100 and arch::Sm120, so a single scale swizzle serves both families and both kernels can cover SM12x. - Parameterise the grouped GEMM on ArchTag through a new Mxfp4GroupGemmArchConfig trait. SM10x/11x keeps its dedicated 1-SM MXFP4 Ptr-Array schedule and fixed epilogue tile unchanged. SM12x has no MXFP4-specific Ptr-Array schedule tag, so it lets the builder derive the cooperative Ptr-Array block-scaled schedule from the pointer-typed StrideA and pick an epilogue tile that fits the tighter SMEM budget of a TMEM-less part. - Widen the experts-quant capability guard to include SM12x. - Compile both MXFP4 sources for SM12x. - Gate the MXFP4 MoE kernel tests on the compiled-capability query rather than a hardcoded SM family, so they exercise every Blackwell variant the wheel was built for instead of silently skipping. Signed-off-by: ayrnb <641876696@qq.com> --- CMakeLists.txt | 14 ++- .../fp4/mxfp4_blockwise_moe_kernel.cu | 100 +++++++++++++----- .../quantization/fp4/mxfp4_experts_quant.cu | 10 +- tests/kernels/moe/test_mxfp4_moe.py | 29 ++--- 4 files changed, 106 insertions(+), 47 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 54374c2d01c0..4ccffc1f6329 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -976,10 +976,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP") 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" + "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu" "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_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") + if(NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9) + message(STATUS + "Building mxfp4_experts_quant unsupported stubs for SM12x because CUDA compiler version is not >= 12.9 (found ${CMAKE_CUDA_COMPILER_VERSION}).") + endif() set_gencode_flags_for_srcs( SRCS "${FP4_SM120_SRCS}" CUDA_ARCHS "${FP4_SM120_ARCHS}") @@ -991,8 +997,10 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP") message(STATUS "Not building SM12x NVFP4 as no compatible archs were found.") endif() - # SM10x/11x FP4 kernels. MXFP4 experts quantization is currently compiled - # only in this block; SM12x has separate NVFP4 matmul/MoE kernels above. + # SM10x/11x FP4 kernels. The MXFP4 experts-quant and grouped-GEMM sources + # are shared with the SM12x block above: CUTLASS resolves the same + # block-scaled scale-factor layout for arch::Sm100 and arch::Sm120, so a + # single SF swizzle serves both families. if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0) cuda_archs_loose_intersection(FP4_SM100_ARCHS "10.0f;10.7f;11.0f" "${CUDA_ARCHS}") else() diff --git a/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu b/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu index e4d2f2201250..bf13187f86f4 100644 --- a/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu +++ b/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu @@ -2,8 +2,9 @@ * SPDX-License-Identifier: Apache-2.0 * SPDX-FileCopyrightText: Copyright contributors to the vLLM project * - * MXFP4 x MXFP4 block-scaled grouped GEMM kernel for MoE on SM100. - * Uses Cutlass mx_float4_t operands, E8M0 block scales, and 32-element groups. + * MXFP4 x MXFP4 block-scaled grouped GEMM kernel for MoE on SM10x/11x and + * SM12x. Uses Cutlass mx_float4_t operands, E8M0 block scales, and + * 32-element groups. */ #include @@ -166,8 +167,41 @@ void mxfp4_run_get_group_gemm_starts( } } -template -void run_mxfp4_blockwise_scaled_group_mm_sm100( +// Architecture-specific kernel configuration. +// +// SM100 (datacenter Blackwell) has a dedicated 1-SM MXFP4 Ptr-Array schedule +// and a 2-SM cluster capable block-scaled MMA. SM120 (GeForce/RTX PRO +// Blackwell) has no TMEM and no 2-SM block-scaled cluster, so it relies on +// the generic cooperative Ptr-Array block-scaled schedule that CUTLASS +// auto-selects from a pointer-typed StrideA. +template +struct Mxfp4GroupGemmArchConfig; + +#if defined ENABLE_NVFP4_SM100 && ENABLE_NVFP4_SM100 +template <> +struct Mxfp4GroupGemmArchConfig { + using ClusterShape = Shape<_1, _1, _1>; + using MmaTileShape = Shape<_128, _128, _128>; + using EpilogueTile = Shape<_128, _64>; + using KernelSchedule = + cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmMxf4Sm100; + using EpilogueSchedule = cutlass::epilogue::PtrArrayTmaWarpSpecialized1Sm; +}; +#endif + +#if defined ENABLE_NVFP4_SM120 && ENABLE_NVFP4_SM120 +template <> +struct Mxfp4GroupGemmArchConfig { + using ClusterShape = Shape<_1, _1, _1>; + using MmaTileShape = Shape<_128, _128, _128>; + using EpilogueTile = cutlass::epilogue::collective::EpilogueTileAuto; + using KernelSchedule = cutlass::gemm::collective::KernelScheduleAuto; + using EpilogueSchedule = cutlass::epilogue::collective::EpilogueScheduleAuto; +}; +#endif + +template +void run_mxfp4_blockwise_scaled_group_mm_impl( torch::stable::Tensor& output, const torch::stable::Tensor& a, const torch::stable::Tensor& b, const torch::stable::Tensor& a_blockscale, const torch::stable::Tensor& b_blockscales, @@ -198,35 +232,31 @@ void run_mxfp4_blockwise_scaled_group_mm_sm100( static constexpr int AlignmentD = 128 / cutlass::sizeof_bits::value; // Architecture definitions - using ArchTag = cutlass::arch::Sm100; + using ArchTag = Arch; using EpilogueOperatorClass = cutlass::arch::OpClassTensorOp; using MainloopOperatorClass = cutlass::arch::OpClassBlockScaledTensorOp; using StageCountType = cutlass::gemm::collective::StageCountAuto; - using ClusterShape = Shape<_1, _1, _1>; - struct MMA1SMConfig { - using MmaTileShape = Shape<_128, _128, _128>; - using KernelSchedule = - cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmMxf4Sm100; - using EpilogueSchedule = cutlass::epilogue::PtrArrayTmaWarpSpecialized1Sm; - }; + using ArchConfig = Mxfp4GroupGemmArchConfig; + using ClusterShape = typename ArchConfig::ClusterShape; + using MmaTileShape = typename ArchConfig::MmaTileShape; using CollectiveEpilogue = - typename cutlass::epilogue::collective::CollectiveBuilder< - ArchTag, EpilogueOperatorClass, typename MMA1SMConfig::MmaTileShape, - ClusterShape, Shape<_128, _64>, ElementAccumulator, + typename cutlass::epilogue::collective::CollectiveBuilder< + ArchTag, EpilogueOperatorClass, MmaTileShape, ClusterShape, + typename ArchConfig::EpilogueTile, ElementAccumulator, ElementAccumulator, ElementC, LayoutC*, AlignmentC, ElementD, LayoutC*, AlignmentD, - typename MMA1SMConfig::EpilogueSchedule>::CollectiveOp; + typename ArchConfig::EpilogueSchedule>::CollectiveOp; using CollectiveMainloop = - typename cutlass::gemm::collective::CollectiveBuilder< + typename cutlass::gemm::collective::CollectiveBuilder< ArchTag, MainloopOperatorClass, ElementA, LayoutA*, AlignmentA, - ElementB, LayoutB*, AlignmentB, ElementAccumulator, - typename MMA1SMConfig::MmaTileShape, ClusterShape, - cutlass::gemm::collective::StageCountAutoCarveout( - sizeof(typename CollectiveEpilogue::SharedStorage))>, - typename MMA1SMConfig::KernelSchedule>::CollectiveOp; + ElementB, LayoutB*, AlignmentB, ElementAccumulator, MmaTileShape, + ClusterShape, + cutlass::gemm::collective::StageCountAutoCarveout( + sizeof(typename CollectiveEpilogue::SharedStorage))>, + typename ArchConfig::KernelSchedule>::CollectiveOp; using GemmKernel = cutlass::gemm::kernel::GemmUniversal= 100 && version_num < 120) { - run_mxfp4_blockwise_scaled_group_mm_sm100( + run_mxfp4_blockwise_scaled_group_mm_impl( output, a, b, a_blockscale, b_blockscales, problem_sizes, expert_offsets, sf_offsets, M, N, K); + return; + } +#endif +#if defined ENABLE_NVFP4_SM120 && ENABLE_NVFP4_SM120 + if (version_num >= 120 && version_num < 130) { + run_mxfp4_blockwise_scaled_group_mm_impl( +output, a, b, a_blockscale, b_blockscales, problem_sizes, + expert_offsets, sf_offsets, M, N, K); return; } #endif STD_TORCH_CHECK_NOT_IMPLEMENTED( false, "No compiled cutlass_mxfp4_group_mm kernel for CUDA device capability: ", - version_num, ". Required capability: 100"); + version_num, ". Required capability: 100-119 (SM10x/11x) or 120-129 " + "(SM12x)."); } -#if defined ENABLE_NVFP4_SM100 && ENABLE_NVFP4_SM100 +#if (defined ENABLE_NVFP4_SM100 && ENABLE_NVFP4_SM100) || \ + (defined ENABLE_NVFP4_SM120 && ENABLE_NVFP4_SM120) constexpr auto MXFP4_FLOAT4_E2M1X2 = torch::headeronly::ScalarType::Byte; // E8M0 scale factors stored as uint8 constexpr auto MXFP4_SF_DTYPE = torch::headeronly::ScalarType::Byte; @@ -417,8 +457,9 @@ void cutlass_mxfp4_group_mm(torch::stable::Tensor& output, const torch::stable::Tensor& problem_sizes, const torch::stable::Tensor& expert_offsets, const torch::stable::Tensor& sf_offsets) { -#if defined ENABLE_NVFP4_SM100 && ENABLE_NVFP4_SM100 - // Input validation +#if (defined ENABLE_NVFP4_SM100 && ENABLE_NVFP4_SM100) || \ + (defined ENABLE_NVFP4_SM120 && ENABLE_NVFP4_SM120) +// Input validation CHECK_INPUT(a, MXFP4_FLOAT4_E2M1X2, "a"); CHECK_INPUT(b, MXFP4_FLOAT4_E2M1X2, "b"); // MXFP4 uses E8M0 scale factors (stored as uint8) @@ -462,8 +503,9 @@ void cutlass_mxfp4_group_mm(torch::stable::Tensor& output, #else STD_TORCH_CHECK_NOT_IMPLEMENTED( false, - "No compiled cutlass_mxfp4_group_mm kernel; build vLLM with " - "SM100 block-scaled FP4 MoE (ENABLE_NVFP4_SM100) and CUDA 12.8+."); + "No compiled cutlass_mxfp4_group_mm kernel; build vLLM with " + "SM10x/11x (ENABLE_NVFP4_SM100) or SM12x (ENABLE_NVFP4_SM120) " + "block-scaled FP4 MoE and CUDA 12.8+."); #endif } diff --git a/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu b/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu index 20f024bcef51..6f3d541788c2 100644 --- a/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu +++ b/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu @@ -377,7 +377,11 @@ static void validate_mxfp4_experts_quant_inputs( static bool mxfp4_experts_quant_sm_supported(int64_t cuda_device_capability) { #if VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED - return cuda_device_capability >= 100 && cuda_device_capability < 120; + // SM10x/11x and SM12x share this kernel: CUTLASS resolves the same + // Sm1xxBlkScaledConfig scale-factor layout for cutlass::arch::Sm100 and + // cutlass::arch::Sm120, so the swizzle emitted by + // cvt_quant_to_fp4_get_sf_out_offset() is valid for both families. + return cuda_device_capability >= 100 && cuda_device_capability < 130; #else return false; #endif @@ -393,7 +397,7 @@ void mxfp4_experts_quant( int32_t sm = get_sm_version_num(); STD_TORCH_CHECK(mxfp4_experts_quant_sm_supported(sm), "No compiled MXFP4 experts quant kernel for SM ", sm, - ". Recompile with SM10x/11x FP4 support and CUDA >= 12.9."); + ". Recompile with SM10x/11x or SM12x FP4 support and CUDA >= 12.9."); auto m_topk = input.size(0); auto k = input.size(1); @@ -431,7 +435,7 @@ void silu_and_mul_mxfp4_experts_quant( int32_t sm = get_sm_version_num(); STD_TORCH_CHECK(mxfp4_experts_quant_sm_supported(sm), "No compiled SiLU+Mul MXFP4 experts quant kernel for SM ", sm, - ". Recompile with SM10x/11x FP4 support and CUDA >= 12.9."); + ". Recompile with SM10x/11x or SM12x FP4 support and CUDA >= 12.9."); auto m_topk = input.size(0); auto k_times_2 = input.size(1); diff --git a/tests/kernels/moe/test_mxfp4_moe.py b/tests/kernels/moe/test_mxfp4_moe.py index 16b233b935e0..63c9d493d10f 100644 --- a/tests/kernels/moe/test_mxfp4_moe.py +++ b/tests/kernels/moe/test_mxfp4_moe.py @@ -1,7 +1,7 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project -"""Tests for SM100 CUTLASS MXFP4 x MXFP4 grouped MoE kernels.""" +"""Tests for CUTLASS MXFP4 x MXFP4 grouped MoE kernels (SM10x/11x and SM12x).""" import random @@ -30,9 +30,14 @@ def calc_diff(x, y): return 1 - sim -def is_sm100_supported() -> bool: - return current_platform.is_cuda() and current_platform.is_device_capability_family( - 100 +def is_mxfp4_moe_supported() -> bool: + """Ask the build which devices carry compiled MXFP4 MoE kernels. + """ + capability = current_platform.get_device_capability() + return ( + current_platform.is_cuda() + and capability is not None + and ops.mxfp4_experts_quant_supported(capability.to_int()) ) @@ -61,8 +66,8 @@ def compute_ref_output( @pytest.mark.skipif( - not is_sm100_supported(), - reason="cutlass_mxfp4_group_mm requires CUDA SM100", + not is_mxfp4_moe_supported(), + reason="cutlass_mxfp4_group_mm requires a build with MXFP4 MoE kernels", ) @pytest.mark.parametrize("num_experts", [8, 16, 32]) @pytest.mark.parametrize("out_dtype", [torch.bfloat16]) @@ -201,8 +206,8 @@ def test_cutlass_mxfp4_grouped_mm(num_experts, out_dtype): @pytest.mark.skipif( - not is_sm100_supported(), - reason="mxfp4_experts_quant requires CUDA SM100", + not is_mxfp4_moe_supported(), + reason="mxfp4_experts_quant requires a build with MXFP4 MoE kernels", ) def test_mxfp4_experts_quant_basic(): """ @@ -289,8 +294,8 @@ def compute_reference_e8m0_scale(block_max: float) -> int: @pytest.mark.skipif( - not is_sm100_supported(), - reason="mxfp4_experts_quant requires CUDA SM100", + not is_mxfp4_moe_supported(), + reason="mxfp4_experts_quant requires a build with MXFP4 MoE kernels", ) @pytest.mark.parametrize("k", [256, 7168]) @pytest.mark.parametrize("m", [16, 64]) @@ -412,8 +417,8 @@ def test_mxfp4_experts_quant_e8m0_scale_correctness(m, k): @pytest.mark.skipif( - not is_sm100_supported(), - reason="mxfp4_experts_quant requires CUDA SM100", + not is_mxfp4_moe_supported(), + reason="mxfp4_experts_quant requires a build with MXFP4 MoE kernels", ) def test_mxfp4_experts_quant_no_saturation(): """ From cb5e1f3a7b0d890c3ebb70b8d8aadc7fa37dad7f Mon Sep 17 00:00:00 2001 From: ayrnb <641876696@qq.com> Date: Thu, 6 Aug 2026 02:57:20 +0000 Subject: [PATCH 2/2] Lint code Signed-off-by: ayrnb <641876696@qq.com> --- CMakeLists.txt | 4 +-- .../fp4/mxfp4_blockwise_moe_kernel.cu | 25 ++++++++++--------- .../quantization/fp4/mxfp4_experts_quant.cu | 14 ++++++----- tests/kernels/moe/test_mxfp4_moe.py | 3 +-- 4 files changed, 24 insertions(+), 22 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 4ccffc1f6329..fea927d75bfb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -976,11 +976,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP") 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" + "csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu" "csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_sm120_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/quantization/fp4/mxfp4_blockwise_moe_kernel.cu" "csrc/libtorch_stable/nvfp4_kv_cache_kernels.cu") if(NOT ${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.9) message(STATUS diff --git a/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu b/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu index bf13187f86f4..58fb3636908a 100644 --- a/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu +++ b/csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu @@ -184,7 +184,7 @@ struct Mxfp4GroupGemmArchConfig { using MmaTileShape = Shape<_128, _128, _128>; using EpilogueTile = Shape<_128, _64>; using KernelSchedule = - cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmMxf4Sm100; + cutlass::gemm::KernelPtrArrayTmaWarpSpecialized1SmMxf4Sm100; using EpilogueSchedule = cutlass::epilogue::PtrArrayTmaWarpSpecialized1Sm; }; #endif @@ -242,7 +242,7 @@ void run_mxfp4_blockwise_scaled_group_mm_impl( using MmaTileShape = typename ArchConfig::MmaTileShape; using CollectiveEpilogue = - typename cutlass::epilogue::collective::CollectiveBuilder< + typename cutlass::epilogue::collective::CollectiveBuilder< ArchTag, EpilogueOperatorClass, MmaTileShape, ClusterShape, typename ArchConfig::EpilogueTile, ElementAccumulator, ElementAccumulator, ElementC, LayoutC*, AlignmentC, ElementD, @@ -250,12 +250,12 @@ void run_mxfp4_blockwise_scaled_group_mm_impl( typename ArchConfig::EpilogueSchedule>::CollectiveOp; using CollectiveMainloop = - typename cutlass::gemm::collective::CollectiveBuilder< + typename cutlass::gemm::collective::CollectiveBuilder< ArchTag, MainloopOperatorClass, ElementA, LayoutA*, AlignmentA, ElementB, LayoutB*, AlignmentB, ElementAccumulator, MmaTileShape, ClusterShape, - cutlass::gemm::collective::StageCountAutoCarveout( - sizeof(typename CollectiveEpilogue::SharedStorage))>, + cutlass::gemm::collective::StageCountAutoCarveout( + sizeof(typename CollectiveEpilogue::SharedStorage))>, typename ArchConfig::KernelSchedule>::CollectiveOp; using GemmKernel = @@ -412,21 +412,22 @@ void run_mxfp4_blockwise_scaled_group_mm( run_mxfp4_blockwise_scaled_group_mm_impl( output, a, b, a_blockscale, b_blockscales, problem_sizes, expert_offsets, sf_offsets, M, N, K); - return; + return; } #endif #if defined ENABLE_NVFP4_SM120 && ENABLE_NVFP4_SM120 if (version_num >= 120 && version_num < 130) { run_mxfp4_blockwise_scaled_group_mm_impl( -output, a, b, a_blockscale, b_blockscales, problem_sizes, - expert_offsets, sf_offsets, M, N, K); + output, a, b, a_blockscale, b_blockscales, problem_sizes, + expert_offsets, sf_offsets, M, N, K); return; } #endif STD_TORCH_CHECK_NOT_IMPLEMENTED( false, "No compiled cutlass_mxfp4_group_mm kernel for CUDA device capability: ", - version_num, ". Required capability: 100-119 (SM10x/11x) or 120-129 " + version_num, + ". Required capability: 100-119 (SM10x/11x) or 120-129 " "(SM12x)."); } @@ -459,7 +460,7 @@ void cutlass_mxfp4_group_mm(torch::stable::Tensor& output, const torch::stable::Tensor& sf_offsets) { #if (defined ENABLE_NVFP4_SM100 && ENABLE_NVFP4_SM100) || \ (defined ENABLE_NVFP4_SM120 && ENABLE_NVFP4_SM120) -// Input validation + // Input validation CHECK_INPUT(a, MXFP4_FLOAT4_E2M1X2, "a"); CHECK_INPUT(b, MXFP4_FLOAT4_E2M1X2, "b"); // MXFP4 uses E8M0 scale factors (stored as uint8) @@ -503,8 +504,8 @@ void cutlass_mxfp4_group_mm(torch::stable::Tensor& output, #else STD_TORCH_CHECK_NOT_IMPLEMENTED( false, - "No compiled cutlass_mxfp4_group_mm kernel; build vLLM with " - "SM10x/11x (ENABLE_NVFP4_SM100) or SM12x (ENABLE_NVFP4_SM120) " + "No compiled cutlass_mxfp4_group_mm kernel; build vLLM with " + "SM10x/11x (ENABLE_NVFP4_SM100) or SM12x (ENABLE_NVFP4_SM120) " "block-scaled FP4 MoE and CUDA 12.8+."); #endif } diff --git a/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu b/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu index 6f3d541788c2..219cd3c98e7f 100644 --- a/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu +++ b/csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu @@ -395,9 +395,10 @@ void mxfp4_experts_quant( int64_t n_experts) { #if VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED int32_t sm = get_sm_version_num(); - STD_TORCH_CHECK(mxfp4_experts_quant_sm_supported(sm), - "No compiled MXFP4 experts quant kernel for SM ", sm, - ". Recompile with SM10x/11x or SM12x FP4 support and CUDA >= 12.9."); + STD_TORCH_CHECK( + mxfp4_experts_quant_sm_supported(sm), + "No compiled MXFP4 experts quant kernel for SM ", sm, + ". Recompile with SM10x/11x or SM12x FP4 support and CUDA >= 12.9."); auto m_topk = input.size(0); auto k = input.size(1); @@ -433,9 +434,10 @@ void silu_and_mul_mxfp4_experts_quant( int64_t n_experts) { #if VLLM_MXFP4_EXPERTS_QUANT_SUPPORTED int32_t sm = get_sm_version_num(); - STD_TORCH_CHECK(mxfp4_experts_quant_sm_supported(sm), - "No compiled SiLU+Mul MXFP4 experts quant kernel for SM ", sm, - ". Recompile with SM10x/11x or SM12x FP4 support and CUDA >= 12.9."); + STD_TORCH_CHECK( + mxfp4_experts_quant_sm_supported(sm), + "No compiled SiLU+Mul MXFP4 experts quant kernel for SM ", sm, + ". Recompile with SM10x/11x or SM12x FP4 support and CUDA >= 12.9."); auto m_topk = input.size(0); auto k_times_2 = input.size(1); diff --git a/tests/kernels/moe/test_mxfp4_moe.py b/tests/kernels/moe/test_mxfp4_moe.py index 63c9d493d10f..2d8192cb94f5 100644 --- a/tests/kernels/moe/test_mxfp4_moe.py +++ b/tests/kernels/moe/test_mxfp4_moe.py @@ -31,8 +31,7 @@ def calc_diff(x, y): def is_mxfp4_moe_supported() -> bool: - """Ask the build which devices carry compiled MXFP4 MoE kernels. - """ + """Ask the build which devices carry compiled MXFP4 MoE kernels.""" capability = current_platform.get_device_capability() return ( current_platform.is_cuda()