sgl-project/sglang · error · NotImplementedError

Block sparsity + sheared bias is not supported on SM90

Error message

Block sparsity + sheared bias is not supported on SM90

What it means

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.

Source

Thrown at python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_sm90.py:338

        )

        tiled_mma_qk, tiled_mma_pv = self._get_tiled_mma()
        self.num_mma_threads = tiled_mma_qk.size
        self.num_threads_per_warp_group = 128
        self.num_wg_mma = self.num_mma_threads // self.num_threads_per_warp_group
        assert self.num_wg_mma in [1, 2, 3]
        self.num_threads = self.num_threads_per_warp_group * (self.num_wg_mma + 1)
        self.num_producer_threads = 32
        self.num_Q_load_threads = self.num_threads_per_warp_group  # If not TMA_Q
        self.num_epilogue_threads = self.num_mma_threads
        self.num_mma_regs, self.num_producer_regs = {
            1: (256, 56),
            2: (240, 24),
            3: (160, 32),
        }[self.num_wg_mma]
        self.use_block_sparsity = cutlass.const_expr(blocksparse_tensors is not None)
        if cutlass.const_expr(self.use_block_sparsity and self.has_bias):
            raise NotImplementedError(
                "Block sparsity + sheared bias is not supported on SM90"
            )

        self.use_scheduler_barrier = (
            (self.num_wg_mma >= 2 and self.tile_hdim <= 128)
            if const_expr(self.intra_wg_overlap)
            else (self.num_wg_mma == 2)
        )
        self.use_tma_Q = self.arch >= Arch.sm_90 and not (
            self.pack_gqa and self.tile_m % self.qhead_per_kvhead != 0
        )
        # Split partials are float32; store them straight from registers (no TMA O).
        self.use_tma_O = self.use_tma_Q and not self.is_split_kv
        # Producer needs more registers when doing cp.async Q or KV loads
        if const_expr(
            self.num_wg_mma == 2 and (not self.use_tma_Q or not self.use_tma_KV)
        ):
            self.num_mma_regs, self.num_producer_regs = 224, 40

View on GitHub (pinned to 0132848349)

Solutions

  1. Remove the bias (drop the bias tensor) to use block sparsity
  2. Or disable block sparsity to keep the bias

Example fix

// before
attn(..., bias=bias, blocksparse_tensors=bs)
// after
attn(..., bias=bias)  # no blocksparse_tensors
Defensive patterns

Strategy: type-guard

Validate before calling

on_sm90 = torch.cuda.get_device_capability()[0] == 9
if on_sm90 and blocksparse_tensors is not None and bias is not None:
    blocksparse_tensors = None  # drop one feature

Type guard

def sm90_bs_bias_ok(arch, blocksparse_tensors, bias) -> bool:
    return not (arch // 10 == 9 and blocksparse_tensors is not None and bias is not None)

Try / catch

try:
    fa(..., bias=bias, blocksparse_tensors=bs)
except NotImplementedError:
    fa(..., bias=bias)

Prevention

When it happens

Trigger: Calling the SM90 kernel with blocksparse_tensors provided while has_bias is true (a bias tensor was passed).

Common situations: Running a block-sparse attention model that also uses attention bias (e.g. learned bias or ALiBi materialized as bias) on H100/H800.

Related errors

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