sgl-project/sglang · error · ValueError

nvfp4_gemm_swiglu_nvfp4_quant requires SM100, got SM{major}{

Error message

nvfp4_gemm_swiglu_nvfp4_quant requires SM100, got SM{major}{minor}

What it means

The NVFP4 fused GEMM+SwiGLU kernel uses SM100 (Blackwell, compute capability 10.x) tcgen05 MMA instructions that do not exist on earlier architectures. The op checks compute capability and refuses to run on anything other than major version 10.

Source

Thrown at python/sglang/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.py:2870

        alpha: GEMM global dequant scale, scalar or ``[1, 1]``.
        output_global_scale: Output quantization scale-up factor (= 1 /
            down_proj.input_scale_inv).
        enable_pdl: Enable Programmatic Dependent Launch for the fused kernel.

    Returns:
        ``(out_fp4, out_scale)`` directly consumable by the NVFP4 ``down_proj``.
    """
    if ab_dtype != "float4_e2m1fn" or c_dtype != "float4_e2m1fn":
        raise ValueError(
            "nvfp4_gemm_swiglu_nvfp4_quant currently supports NVFP4 input "
            "and output only"
        )
    if a.device.type != "cuda" or b.device.type != "cuda":
        raise ValueError("nvfp4_gemm_swiglu_nvfp4_quant requires CUDA tensors")

    major, minor = get_compute_capability(a.device)
    if major != 10:
        raise ValueError(
            f"nvfp4_gemm_swiglu_nvfp4_quant requires SM100, got SM{major}{minor}"
        )

    m = a.shape[0]
    k = a.shape[1] * 2
    n = b.shape[0]
    if b.shape[1] * 2 != k:
        raise ValueError(f"Shape mismatch: A K={k}, B K={b.shape[1] * 2}")
    if n % 2 != 0:
        raise ValueError(f"Interleaved FC1 N must be even, got {n}")

    l = 1
    n_out = n // 2
    if n_out % sf_vec_size != 0:
        raise ValueError(
            f"Output N={n_out} must be divisible by sf_vec_size={sf_vec_size}"
        )
    scale_n_out = n_out // sf_vec_size

View on GitHub (pinned to 0132848349)

Solutions

Run on a B200/GB200/RTX Blackwell (SM100) GPU
Otherwise disable the fused NVFP4 path and select a quant config supported by your architecture (e.g. FP8 on Hopper)

Example fix

# before
--quantization nvfp4  (on H100)
# after
--quantization fp8     # or run on a Blackwell SM100 GPU

Defensive patterns

Strategy: validation

Validate before calling

major, _ = get_compute_capability(a.device)
assert major == 10, f'requires SM100, got SM{major}'

Type guard

def supports_nvfp4_fused(device):
    return device.type == 'cuda' and torch.cuda.get_device_capability(device)[0] == 10

Prevention

Gate fused NVFP4 paths on compute capability at engine init
Provide an arch-based quant-config fallback (fp8 on SM90)

When it happens

Trigger: Running on Hopper (SM90), Ada (SM89), Ampere (SM80) or older GPUs; e.g. H100 returns major=9 and triggers this error.

Common situations: Deploying a Blackwell-targeted quant config (NVFP4 fused path) on an H100/A100 cluster, or CI runners with older GPUs picking up the fused path by default.

Related errors

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/3419ac2f4c90d9ec. Report an issue: GitHub.