PR #35568: [Bugfix] Fix SM121 (DGX Spark) exclusion from Marlin/CUTLASS FP8 paths

• Repository: vllm-project/vllm
• Author: blake-snc
• State: open | merged: False
• Link: https://github.com/vllm-project/vllm/pull/35568

Description (problem / solution / changelog)

Summary

SM121 (DGX Spark GB10) shares the same FP8 MMA capabilities as SM120 (RTX 5090) — both support native mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32. However, SM121 is excluded from all Marlin and CUTLASS FP8 codepaths by exact-match arch guards (== 120, in [89, 120], enable_sm120_only).

This fixes 8 locations across codegen, runtime, dispatch, and tests using bounded SM12x family checks (arch // 10 == 12, major_capability == 12, enable_sm120_family, is_device_capability_family(120)):

Codegen (FP8 kernel template generation):

• csrc/quantization/marlin/generate_kernels.py: arch in [89, 120] → arch == 89 or arch // 10 == 12
• csrc/moe/marlin_moe_wna16/generate_kernels.py: same fix

Runtime (FP8 activation gate):

• csrc/moe/marlin_moe_wna16/ops.cu: == 120 → major_capability == 12

CUTLASS FP8 dispatch (kernel wrapper):

• csrc/quantization/w8a8/cutlass/c3x/scaled_mm.cuh: enable_sm120_only → enable_sm120_family
• csrc/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8_dispatch.cuh: same fix

Tests (FP8 test case generation):

• tests/kernels/moe/test_moe.py: get_device_capability() not in [89, 120] → proper is_device_capability(89) / is_device_capability_family(120) API calls
• tests/kernels/quantization/test_marlin_gemm.py: same fix

Python-side FP8 input validation:

• vllm/model_executor/layers/quantization/utils/marlin_utils.py: is_device_capability(120) → is_device_capability_family(120)

All checks use bounded SM12x family matching (covers SM120/SM121 but won't accidentally match future SM13x).

The enable_sm120_only → enable_sm120_family change in the CUTLASS dispatch headers also resolves the CUTLASS FP4 GEMM failure on SM121 reported in #30163 ("Failed to run cutlass FP4 gemm on sm120. Error: Error Internal"), since enable_sm120_only uses __CUDA_ARCH__ == 1200 which excludes SM121 (__CUDA_ARCH__ == 1210), while enable_sm120_family uses >= 1200 && < 1300.

Validation

Tested on DGX Spark (NVIDIA GB10, SM121a / capability 12.1):

Marlin FP4 GEMM (all 5 configs including N=100544): PASS CUTLASS FP4 dispatch: cutlass_scaled_mm_supports_fp4(121) = True Capability check logic:

SM89 (Ada):   allowed via exact match ✓
SM90 (Hopper): blocked ✓
SM120 (RTX 5090): allowed ✓
SM121 (DGX Spark): allowed ✓
SM130 (future): not matched ✓

Subsumes #35803. Fixes #35432. Fixes #30163. Relates to #30135.

Contributed by Second Nature Computing (https://joinsecondnature.com)

Test plan

• Validated on SM121a hardware (DGX Spark)
• Marlin FP4 GEMM passes all 5 test configs
• enable_sm120_family verified in common.hpp with correct >= 1200 && < 1300 range guard
• is_device_capability_family(120) verified: uses to_int() // 10 == 120 // 10
• Pre-commit hooks pass

🤖 Generated with Claude Code

Changed files

• csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm.cuh (modified, +1/-1)
• csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8_dispatch.cuh (modified, +1/-1)
• csrc/moe/marlin_moe_wna16/generate_kernels.py (modified, +3/-3)
• csrc/moe/marlin_moe_wna16/ops.cu (modified, +2/-2)
• csrc/quantization/marlin/generate_kernels.py (modified, +3/-3)
• tests/kernels/moe/test_moe.py (modified, +5/-1)
• tests/kernels/quantization/test_marlin_gemm.py (modified, +2/-1)
• vllm/model_executor/layers/quantization/utils/marlin_utils.py (modified, +2/-2)