{"record":{"id":"3419ac2f4c90d9ec","repo":"sgl-project/sglang","slug":"nvfp4-gemm-swiglu-nvfp4-quant-requires-sm100-got","errorCode":null,"errorMessage":"nvfp4_gemm_swiglu_nvfp4_quant requires SM100, got SM{major}{minor}","messagePattern":"nvfp4_gemm_swiglu_nvfp4_quant requires SM100, got SM(.+?)(.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.py","lineNumber":2870,"sourceCode":"        alpha: GEMM global dequant scale, scalar or ``[1, 1]``.\n        output_global_scale: Output quantization scale-up factor (= 1 /\n            down_proj.input_scale_inv).\n        enable_pdl: Enable Programmatic Dependent Launch for the fused kernel.\n\n    Returns:\n        ``(out_fp4, out_scale)`` directly consumable by the NVFP4 ``down_proj``.\n    \"\"\"\n    if ab_dtype != \"float4_e2m1fn\" or c_dtype != \"float4_e2m1fn\":\n        raise ValueError(\n            \"nvfp4_gemm_swiglu_nvfp4_quant currently supports NVFP4 input \"\n            \"and output only\"\n        )\n    if a.device.type != \"cuda\" or b.device.type != \"cuda\":\n        raise ValueError(\"nvfp4_gemm_swiglu_nvfp4_quant requires CUDA tensors\")\n\n    major, minor = get_compute_capability(a.device)\n    if major != 10:\n        raise ValueError(\n            f\"nvfp4_gemm_swiglu_nvfp4_quant requires SM100, got SM{major}{minor}\"\n        )\n\n    m = a.shape[0]\n    k = a.shape[1] * 2\n    n = b.shape[0]\n    if b.shape[1] * 2 != k:\n        raise ValueError(f\"Shape mismatch: A K={k}, B K={b.shape[1] * 2}\")\n    if n % 2 != 0:\n        raise ValueError(f\"Interleaved FC1 N must be even, got {n}\")\n\n    l = 1\n    n_out = n // 2\n    if n_out % sf_vec_size != 0:\n        raise ValueError(\n            f\"Output N={n_out} must be divisible by sf_vec_size={sf_vec_size}\"\n        )\n    scale_n_out = n_out // sf_vec_size","sourceCodeStart":2852,"sourceCodeEnd":2888,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.py#L2852-L2888","documentation":"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.","triggerScenarios":"Running on Hopper (SM90), Ada (SM89), Ampere (SM80) or older GPUs; e.g. H100 returns major=9 and triggers this error.","commonSituations":"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.","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)"],"exampleFix":"# before\n--quantization nvfp4  (on H100)\n# after\n--quantization fp8     # or run on a Blackwell SM100 GPU","handlingStrategy":"validation","validationCode":"major, _ = get_compute_capability(a.device)\nassert major == 10, f'requires SM100, got SM{major}'","typeGuard":"def supports_nvfp4_fused(device):\n    return device.type == 'cuda' and torch.cuda.get_device_capability(device)[0] == 10","tryCatchPattern":null,"preventionTips":["Gate fused NVFP4 paths on compute capability at engine init","Provide an arch-based quant-config fallback (fp8 on SM90)"],"tags":["nvfp4","sm100","blackwell","gpu-architecture","cuda"],"backgroundTag":"gpu-compute-capability-unsupported","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}