{"record":{"id":"273665027653bc0a","repo":"sgl-project/sglang","slug":"dsv3-fused-a-gemm-requires-sm90-hopper-or-later","errorCode":null,"errorMessage":"dsv3_fused_a_gemm requires SM90 (Hopper) or later","messagePattern":"dsv3_fused_a_gemm requires SM90 \\(Hopper\\) or later","errorType":"error_code","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/gemm/cutedsl_dsv3_fused_a_gemm.py","lineNumber":310,"sourceCode":"        stream=stream,\n    )\n\n\n_compiled: dict[tuple[int, int, int], object] = {}\n\n\ndef _pick_nstage(num_kt: int, tile_n: int) -> int:\n    nstage = (get_smem_capacity_in_bytes() // 4 - _BAR_I32) // _stage_i32(tile_n)\n    return min(nstage, MAX_NSTAGE, num_kt)\n\n\ndef _pick_tile_n(num_tokens: int) -> int:\n    return 8 if num_tokens <= 8 else 16\n\n\ndef _compiled_kernel(num_kt: int, gemm_m: int, tile_n: int):\n    if get_device_sm() < 90:\n        raise RuntimeError(\"dsv3_fused_a_gemm requires SM90 (Hopper) or later\")\n    if (num_kt, gemm_m, tile_n) not in _compiled:\n        nstage = _pick_nstage(num_kt, tile_n)\n        smem_bytes = (_BAR_I32 + nstage * _stage_i32(tile_n)) * 4\n        k = num_kt * TILE_K\n        w = torch.empty(gemm_m, k, dtype=torch.bfloat16, device=\"cuda\")\n        a = torch.empty(16, k, dtype=torch.bfloat16, device=\"cuda\")\n        o = torch.empty(16, gemm_m, dtype=torch.bfloat16, device=\"cuda\")\n        stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream)\n        _compiled[(num_kt, gemm_m, tile_n)] = cute.compile(\n            _dsv3_fused_a_gemm_host,\n            from_dlpack(w.view(torch.int32)),\n            from_dlpack(a.view(torch.int32)),\n            from_dlpack(o),\n            cutlass.Int32(16),\n            stream,\n            num_kt,\n            gemm_m,\n            smem_bytes,","sourceCodeStart":292,"sourceCodeEnd":328,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/gemm/cutedsl_dsv3_fused_a_gemm.py#L292-L328","documentation":"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.","triggerScenarios":"Calling _dsv3_fused_a_gemm_run (which invokes the cached _compiled_kernel) on Ampere or older GPUs (A100 is SM80).","commonSituations":"Running DeepSeek-V3 MoE attention-fused paths on A100 or consumer pre-Hopper cards, or in CI without GPU-specific dispatch.","solutions":["Run on H100/H200/B200 or later.","Fall back to a standard torch.matmul path when get_device_sm() < 90.","Gate model-level fused-kernel flags by detected SM version."],"exampleFix":"// before\nout = _dsv3_fused_a_gemm_run(a, w)\n// after\nfrom sglang.kernels.utils import get_device_sm\nout = _dsv3_fused_a_gemm_run(a, w) if get_device_sm() >= 90 else torch.matmul(a, w.t())","handlingStrategy":"fallback","validationCode":"if get_device_sm() < 90:\n    out = torch.matmul(a, w.t())\nelse:\n    out = _dsv3_fused_a_gemm_run(a, w)","typeGuard":"def sm90_plus() -> bool:\n    return torch.cuda.get_device_capability(0)[0] >= 9","tryCatchPattern":null,"preventionTips":["Feature-gate fused kernels by SM version.","Test on the minimum target GPU in CI."],"tags":["gemm","dsv3","sm90","hopper","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"}