{"record":{"id":"6ed6077b37385a39","repo":"sgl-project/sglang","slug":"cutedsl-bf16-gemm-requires-an-sm10x-gpu","errorCode":null,"errorMessage":"cutedsl_bf16_gemm requires an SM10x GPU","messagePattern":"cutedsl_bf16_gemm requires an SM10x GPU","errorType":"error_code","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py","lineNumber":1396,"sourceCode":"        return m <= 64\n    if k < 4096:\n        return k >= 3072 and 25 <= m <= 48 and ragged\n    if n <= 6144:\n        if m <= 64:\n            return m <= 32 or k >= 6144 or ragged\n        return m <= 72 and k >= 6144 and ragged\n    if n <= 8192:\n        return m <= 48 or (m <= 63 and ragged)\n    return n <= 12288 and k >= 6144 and (m <= 32 or (m <= 48 and ragged))\n\n\ndef _tgv_bf16_gemm_run(\n    x: torch.Tensor,\n    weight: torch.Tensor,\n    bias: Optional[torch.Tensor],\n) -> torch.Tensor:\n    if not is_sm100_supported():\n        raise RuntimeError(\"cutedsl_bf16_gemm requires an SM10x GPU\")\n    assert x.dtype == torch.bfloat16 and weight.dtype == torch.bfloat16\n    assert x.stride(-1) == 1, \"x must be K-major [M, K]\"\n    assert weight.stride(-1) == 1, \"weight must be K-major [N, K]\"\n    out = torch.empty(\n        (x.shape[0], weight.shape[0]), dtype=torch.bfloat16, device=x.device\n    )\n    if x.shape[0] == 0:\n        # Match cuBLAS/F.linear semantics for empty batches; a 0-CTA launch\n        # would fail with CUDA_ERROR_INVALID_VALUE.\n        return out\n    return _run_tgv(\n        x,\n        weight.t(),\n        bias,\n        out,\n        pdl=True,\n        tactic=_pick_tactic(x.shape[0], weight.shape[0], weight.shape[1]),\n    )","sourceCodeStart":1378,"sourceCodeEnd":1414,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/gemm/cutedsl_bf16_gemm.py#L1378-L1414","documentation":"The cuteDSL TGV bf16 GEMM uses tcgen05 MMA instructions that only exist on SM100+ (Blackwell) GPUs; is_sm100_supported() gates the run function.","triggerScenarios":"Calling _tgv_bf16_gemm_run on Hopper (H100, SM90), Ada, Ampere, or any non-Blackwell GPU.","commonSituations":"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.","solutions":["Run on an SM100+ GPU (B200, GB200, RTX Blackwell).","Let the dispatcher choose a non-TGV backend on older GPUs instead of forcing this path.","Gate calls with is_sm100_supported() in your own code."],"exampleFix":"// before\nout = _tgv_bf16_gemm_run(x, w, b)  # on H100\n// after\nif is_sm100_supported():\n    out = _tgv_bf16_gemm_run(x, w, b)\nelse:\n    out = torch.nn.functional.linear(x, w, b)","handlingStrategy":"fallback","validationCode":"from sglang.kernels.utils import is_sm100_supported\nif not is_sm100_supported():\n    out = torch.nn.functional.linear(x, w, b)  # fallback","typeGuard":"def sm100_ok() -> bool:\n    return torch.cuda.get_device_capability(0)[0] >= 10","tryCatchPattern":null,"preventionTips":["Gate fused GEMM paths on SM version at startup.","Keep a cuBLAS fallback in the dispatcher."],"tags":["gemm","sm100","blackwell","gpu-architecture"],"backgroundTag":"gpu-architecture-not-supported","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}