sgl-project/sglang · error · RuntimeError

dsv3_fused_a_gemm requires SM90 (Hopper) or later

Error message

dsv3_fused_a_gemm requires SM90 (Hopper) or later

What it means

The DeepSeek-V3 fused A-GEMM cuteDSL kernel relies on Hopper (SM90) TMA/wgmma features; get_device_sm() < 90 raises this at compile time. This covers H100+ including Blackwell.

Source

Thrown at python/sglang/kernels/ops/gemm/cutedsl_dsv3_fused_a_gemm.py:310

        stream=stream,
    )


_compiled: dict[tuple[int, int, int], object] = {}


def _pick_nstage(num_kt: int, tile_n: int) -> int:
    nstage = (get_smem_capacity_in_bytes() // 4 - _BAR_I32) // _stage_i32(tile_n)
    return min(nstage, MAX_NSTAGE, num_kt)


def _pick_tile_n(num_tokens: int) -> int:
    return 8 if num_tokens <= 8 else 16


def _compiled_kernel(num_kt: int, gemm_m: int, tile_n: int):
    if get_device_sm() < 90:
        raise RuntimeError("dsv3_fused_a_gemm requires SM90 (Hopper) or later")
    if (num_kt, gemm_m, tile_n) not in _compiled:
        nstage = _pick_nstage(num_kt, tile_n)
        smem_bytes = (_BAR_I32 + nstage * _stage_i32(tile_n)) * 4
        k = num_kt * TILE_K
        w = torch.empty(gemm_m, k, dtype=torch.bfloat16, device="cuda")
        a = torch.empty(16, k, dtype=torch.bfloat16, device="cuda")
        o = torch.empty(16, gemm_m, dtype=torch.bfloat16, device="cuda")
        stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream)
        _compiled[(num_kt, gemm_m, tile_n)] = cute.compile(
            _dsv3_fused_a_gemm_host,
            from_dlpack(w.view(torch.int32)),
            from_dlpack(a.view(torch.int32)),
            from_dlpack(o),
            cutlass.Int32(16),
            stream,
            num_kt,
            gemm_m,
            smem_bytes,

View on GitHub (pinned to 0132848349)

Solutions

1. Run on H100/H200/B200 or later.
2. Fall back to a standard torch.matmul path when get_device_sm() < 90.
3. Gate model-level fused-kernel flags by detected SM version.

Example fix

// before
out = _dsv3_fused_a_gemm_run(a, w)
// after
from sglang.kernels.utils import get_device_sm
out = _dsv3_fused_a_gemm_run(a, w) if get_device_sm() >= 90 else torch.matmul(a, w.t())

if get_device_sm() < 90:
    out = torch.matmul(a, w.t())
else:
    out = _dsv3_fused_a_gemm_run(a, w)

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

When it happens

Trigger: Calling _dsv3_fused_a_gemm_run (which invokes the cached _compiled_kernel) on Ampere or older GPUs (A100 is SM80).

Common situations: Running DeepSeek-V3 MoE attention-fused paths on A100 or consumer pre-Hopper cards, or in CI without GPU-specific dispatch.

Related errors

• qprep_bf16_fp8_sm90 requires an SM90 (Hopper) GPU
• cutedsl_bf16_gemm requires an SM10x GPU
• fp8_blockwise_scaled_mm JIT kernel requires SM120 (Blackwell
• TRTLLM MHA backend for prefill requires Hopper (SM90), Black
• flash_attn at sgl-kernel is only supported on sm90 and above

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