sgl-project/sglang · error · RuntimeError

unsupported input for Sana fused bias-SiLU

Error message

unsupported input for Sana fused bias-SiLU

What it means

fused_bias_silu is a bit-exact Triton replacement for Sana conv bias+SiLU and only accepts x that is a 4-D, non-empty, CUDA bfloat16 tensor contiguous in channels_last memory format, plus a 1-D contiguous bf16 bias on the same device with bias.shape[0] == x.shape[1]. Any deviation raises RuntimeError('unsupported input for Sana fused bias-SiLU').

Source

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

        and x.is_contiguous(memory_format=torch.channels_last)
    )


def can_use_fused_bias_silu(x: torch.Tensor, bias: torch.Tensor) -> bool:
    return (
        _is_channels_last_bf16(x)
        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_silu(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:
    if not can_use_fused_bias_silu(x, bias):
        raise RuntimeError("unsupported input for Sana fused bias-SiLU")
    out = torch.empty_like(x, memory_format=torch.preserve_format)
    with torch.cuda.device(x.device):
        _bias_silu_kernel[(triton.cdiv(x.numel(), 1024),)](
            out, x, bias, x.numel(), channels=x.shape[1]
        )
    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()

View on GitHub (pinned to 0132848349)

Solutions

1. Convert layout: x = x.to(memory_format=torch.channels_last, dtype=torch.bfloat16)
2. Match bias: bias = bias.to(x.dtype, x.device); assert bias.shape[0] == x.shape[1]
3. Guard with can_use_fused_bias_silu(x, bias) and fall back to F.silu(x + bias[:, None, None]) otherwise

Example fix

// before
y = fused_bias_silu(x, bias)  # x is channels_first fp32
// after
if can_use_fused_bias_silu(x, bias):
    y = fused_bias_silu(x, bias)
else:
    y = F.silu(x.float() + bias[:, None, None].float()).to(x.dtype)

from sglang.kernels.ops.diffusion.activation.sana_conv_post_triton import can_use_fused_bias_silu
if not can_use_fused_bias_silu(x, bias):
    x = x.to(memory_format=torch.channels_last, dtype=torch.bfloat16)
    bias = bias.to(torch.bfloat16, x.device)

def usable_bias_silu(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 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_silu(x, bias)
except RuntimeError:
    y = torch.nn.functional.silu(x.float() + bias[:, None, None].float()).to(x.dtype)

When it happens

Trigger: Passing a channels-first contiguous tensor (the PyTorch default after most ops), an fp16 or 3-D/5-D tensor, an empty tensor, or a bias with length != x.shape[1] or on a different device/dtype.

Common situations: A conv layer whose weight/layout config does not produce channels_last output, bf16-vs-fp16 model variants, or CPU-side unit tests calling the fused op without CUDA tensors.

Related errors

• unsupported input for Sana fused bias-GLU
• unsupported input for wan_rmsnorm_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

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