sgl-project/sglang · error · RuntimeError

unsupported input for Sana fused bias-GLU

Error message

unsupported input for Sana fused bias-GLU

What it means

fused_bias_glu requires x to be a 4-D non-empty CUDA bfloat16 tensor contiguous in channels_last with an even channel count (x.shape[1] % 2 == 0, holding hidden+gate halves), and a 1-D contiguous bf16 bias of length x.shape[1] on the same device. Any other contract raises RuntimeError('unsupported input for Sana fused bias-GLU').

Source

Thrown at python/sglang/kernels/ops/diffusion/activation/sana_conv_post_triton.py:102

    return out


def can_use_fused_bias_glu(x: torch.Tensor, bias: torch.Tensor) -> bool:
    return (
        _is_channels_last_bf16(x)
        and x.shape[1] % 2 == 0
        and bias.is_cuda
        and bias.dtype is x.dtype
        and bias.device == x.device
        and bias.dim() == 1
        and bias.shape[0] == x.shape[1]
        and bias.is_contiguous()
    )


def fused_bias_glu(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
    if not can_use_fused_bias_glu(x, bias):
        raise RuntimeError("unsupported input for Sana fused bias-GLU")
    batch, double_channels, height, width = x.shape
    channels = double_channels // 2
    out = torch.empty(
        (batch, channels, height, width),
        dtype=x.dtype,
        device=x.device,
        memory_format=torch.channels_last,
    )
    with torch.cuda.device(x.device):
        _bias_glu_kernel[(triton.cdiv(out.numel(), 1024),)](
            out, x, bias, out.numel(), channels=channels
        )
    return out


__all__ = [
    "can_use_fused_bias_glu",
    "can_use_fused_bias_silu",

View on GitHub (pinned to 0132848349)

Solutions

1. Ensure channels are even and tensor is channels_last bf16: x = x.contiguous(memory_format=torch.channels_last).bfloat16()
2. Verify bias.shape == (x.shape[1],) and bias is on x.device with bf16 dtype
3. Guard with can_use_fused_bias_glu(x, bias) and use an eager fallback

Example fix

// before
y = fused_bias_glu(x, bias)
// after
if can_use_fused_bias_glu(x, bias):
    y = fused_bias_glu(x, bias)
else:
    h, g = x[:, :C].float(), x[:, C:].float()
    y = ((h + bias[:C, None, None]).float() * torch.sigmoid(g + bias[C:, None, None]).float()).to(x.dtype)

from sglang.kernels.ops.diffusion.activation.sana_conv_post_triton import can_use_fused_bias_glu
if not can_use_fused_bias_glu(x, bias):
    raise ValueError('fall back to eager GLU')  # or route to eager path

def usable_bias_glu(x, b) -> bool:
    return (x.is_cuda and x.dtype is torch.bfloat16 and x.dim() == 4
            and x.numel() > 0 and x.is_contiguous(memory_format=torch.channels_last)
            and x.shape[1] % 2 == 0
            and b.is_cuda and b.dtype is x.dtype and b.device == x.device
            and b.dim() == 1 and b.shape[0] == x.shape[1] and b.is_contiguous())

try:
    y = fused_bias_glu(x, bias)
except RuntimeError:
    C = x.shape[1] // 2
    y = eager_bias_glu(x, bias, C)

When it happens

Trigger: Passing a tensor with odd channel count, channels-first layout, fp16/fp32 dtype, or a bias whose length doesn't equal x.shape[1]; also non-CPU-safe calls (CPU tensors, mismatched devices).

Common situations: Wiring a non-Sana conv projection (odd channels) into the Sana GLU path, or forgetting .to(torch.channels_last) after a checkpoint load that resets strides.

Related errors

• unsupported input for Sana fused bias-SiLU
• Unsupported content type ${header.content_type}
• {self._op_label()}: no triton backend
• sparse_attn_v4_paged_decode expects fp16/bf16 q, got {q.dtyp
• kv_scales supplied but unified_kv is {unified_kv.dtype}, exp

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