{"record":{"id":"c25366477c7a2efd","repo":"sgl-project/sglang","slug":"block-sparsity-sheared-bias-is-not-supported-on","errorCode":null,"errorMessage":"Block sparsity + sheared bias is not supported on SM90","messagePattern":"Block sparsity \\+ sheared bias is not supported on SM90","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_sm90.py","lineNumber":338,"sourceCode":"        )\n\n        tiled_mma_qk, tiled_mma_pv = self._get_tiled_mma()\n        self.num_mma_threads = tiled_mma_qk.size\n        self.num_threads_per_warp_group = 128\n        self.num_wg_mma = self.num_mma_threads // self.num_threads_per_warp_group\n        assert self.num_wg_mma in [1, 2, 3]\n        self.num_threads = self.num_threads_per_warp_group * (self.num_wg_mma + 1)\n        self.num_producer_threads = 32\n        self.num_Q_load_threads = self.num_threads_per_warp_group  # If not TMA_Q\n        self.num_epilogue_threads = self.num_mma_threads\n        self.num_mma_regs, self.num_producer_regs = {\n            1: (256, 56),\n            2: (240, 24),\n            3: (160, 32),\n        }[self.num_wg_mma]\n        self.use_block_sparsity = cutlass.const_expr(blocksparse_tensors is not None)\n        if cutlass.const_expr(self.use_block_sparsity and self.has_bias):\n            raise NotImplementedError(\n                \"Block sparsity + sheared bias is not supported on SM90\"\n            )\n\n        self.use_scheduler_barrier = (\n            (self.num_wg_mma >= 2 and self.tile_hdim <= 128)\n            if const_expr(self.intra_wg_overlap)\n            else (self.num_wg_mma == 2)\n        )\n        self.use_tma_Q = self.arch >= Arch.sm_90 and not (\n            self.pack_gqa and self.tile_m % self.qhead_per_kvhead != 0\n        )\n        # Split partials are float32; store them straight from registers (no TMA O).\n        self.use_tma_O = self.use_tma_Q and not self.is_split_kv\n        # Producer needs more registers when doing cp.async Q or KV loads\n        if const_expr(\n            self.num_wg_mma == 2 and (not self.use_tma_Q or not self.use_tma_KV)\n        ):\n            self.num_mma_regs, self.num_producer_regs = 224, 40","sourceCodeStart":320,"sourceCodeEnd":356,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_sm90.py#L320-L356","documentation":"The SM90 (Hopper) FA4 kernel does not support combining block-sparse attention with sheared/attention bias. The tile scheduler cannot apply both sparsity masks and bias tiles simultaneously.","triggerScenarios":"Calling the SM90 kernel with blocksparse_tensors provided while has_bias is true (a bias tensor was passed).","commonSituations":"Running a block-sparse attention model that also uses attention bias (e.g. learned bias or ALiBi materialized as bias) on H100/H800.","solutions":["Remove the bias (drop the bias tensor) to use block sparsity","Or disable block sparsity to keep the bias"],"exampleFix":"// before\nattn(..., bias=bias, blocksparse_tensors=bs)\n// after\nattn(..., bias=bias)  # no blocksparse_tensors","handlingStrategy":"type-guard","validationCode":"on_sm90 = torch.cuda.get_device_capability()[0] == 9\nif on_sm90 and blocksparse_tensors is not None and bias is not None:\n    blocksparse_tensors = None  # drop one feature","typeGuard":"def sm90_bs_bias_ok(arch, blocksparse_tensors, bias) -> bool:\n    return not (arch // 10 == 9 and blocksparse_tensors is not None and bias is not None)","tryCatchPattern":"try:\n    fa(..., bias=bias, blocksparse_tensors=bs)\nexcept NotImplementedError:\n    fa(..., bias=bias)","preventionTips":["Keep a feature matrix of kernel capabilities per SM architecture","Never combine experimental features (blocksparse + bias) without checking support"],"tags":["flash-attention","block-sparse","bias","sm90","unsupported-feature"],"backgroundTag":"unsupported-feature-combination","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}