sgl-project/sglang · error · RuntimeError

cutedsl_bf16_gemm requires an SM10x GPU

Error message

cutedsl_bf16_gemm requires an SM10x GPU

What it means

The cuteDSL TGV bf16 GEMM uses tcgen05 MMA instructions that only exist on SM100+ (Blackwell) GPUs; is_sm100_supported() gates the run function.

Source

Thrown at python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py:1396

        return m <= 64
    if k < 4096:
        return k >= 3072 and 25 <= m <= 48 and ragged
    if n <= 6144:
        if m <= 64:
            return m <= 32 or k >= 6144 or ragged
        return m <= 72 and k >= 6144 and ragged
    if n <= 8192:
        return m <= 48 or (m <= 63 and ragged)
    return n <= 12288 and k >= 6144 and (m <= 32 or (m <= 48 and ragged))


def _tgv_bf16_gemm_run(
    x: torch.Tensor,
    weight: torch.Tensor,
    bias: Optional[torch.Tensor],
) -> torch.Tensor:
    if not is_sm100_supported():
        raise RuntimeError("cutedsl_bf16_gemm requires an SM10x GPU")
    assert x.dtype == torch.bfloat16 and weight.dtype == torch.bfloat16
    assert x.stride(-1) == 1, "x must be K-major [M, K]"
    assert weight.stride(-1) == 1, "weight must be K-major [N, K]"
    out = torch.empty(
        (x.shape[0], weight.shape[0]), dtype=torch.bfloat16, device=x.device
    )
    if x.shape[0] == 0:
        # Match cuBLAS/F.linear semantics for empty batches; a 0-CTA launch
        # would fail with CUDA_ERROR_INVALID_VALUE.
        return out
    return _run_tgv(
        x,
        weight.t(),
        bias,
        out,
        pdl=True,
        tactic=_pick_tactic(x.shape[0], weight.shape[0], weight.shape[1]),
    )

View on GitHub (pinned to 0132848349)

Solutions

1. Run on an SM100+ GPU (B200, GB200, RTX Blackwell).
2. Let the dispatcher choose a non-TGV backend on older GPUs instead of forcing this path.
3. Gate calls with is_sm100_supported() in your own code.

Example fix

// before
out = _tgv_bf16_gemm_run(x, w, b)  # on H100
// after
if is_sm100_supported():
    out = _tgv_bf16_gemm_run(x, w, b)
else:
    out = torch.nn.functional.linear(x, w, b)

from sglang.kernels.utils import is_sm100_supported
if not is_sm100_supported():
    out = torch.nn.functional.linear(x, w, b)  # fallback

def sm100_ok() -> bool:
    return torch.cuda.get_device_capability(0)[0] >= 10

When it happens

Trigger: Calling _tgv_bf16_gemm_run on Hopper (H100, SM90), Ada, Ampere, or any non-Blackwell GPU.

Common situations: Running a Blackwell-optimized build on an older cluster, or CI machines without B200/GB200. The custom-op dispatch usually falls back to cuBLAS on unsupported hardware if configured so.

Related errors

• fp8_blockwise_scaled_mm JIT kernel requires SM120 (Blackwell
• nvfp4_gemm_swiglu_nvfp4_quant requires SM100, got SM{major}{
• cta_n={cta_n} invalid for use_2cta={use_2cta}: bf16 K-major
• dsv3_fused_a_gemm requires SM90 (Hopper) or later
• --kv-cache-dtype mxfp8 requires an SM100+ (Blackwell) GPU fo

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