sgl-project/sglang · error · RuntimeError

unsupported input for packed fused SiLU-mul

Error message

unsupported input for packed fused SiLU-mul

What it means

fused_packed_silu_mul_bitexact requires a CUDA tensor whose last dim is even (packed hidden+gate), last-dim stride is 1, second-to-last stride >= shape[-1], the batch stride relationship x.stride(0) == x.shape[1]*x.stride(1) holds (uniform row grouping across leading dims), and numel > 0. Violations raise RuntimeError('unsupported input for packed fused SiLU-mul').

Source

Thrown at python/sglang/kernels/ops/diffusion/activation/silu_mul_bitexact.py:113

            numel,
            BLOCK=1024,
        )
    return out


def fused_packed_silu_mul_bitexact(x: torch.Tensor) -> torch.Tensor:
    """Bit-exact SwiGLU over a contiguous packed ``[..., 2 * D]`` input."""
    if not (
        x.is_cuda
        and x.dtype is torch.bfloat16
        and x.dim() == 3
        and x.stride(-1) == 1
        and x.stride(-2) >= x.shape[-1]
        and x.stride(0) == x.shape[1] * x.stride(1)
        and x.shape[-1] % 2 == 0
        and x.numel() > 0
    ):
        raise RuntimeError("unsupported input for packed fused SiLU-mul")
    hidden = x.shape[-1] // 2
    rows = x.numel() // x.shape[-1]
    row_stride = x.stride(-2)
    out = torch.empty((*x.shape[:-1], hidden), dtype=x.dtype, device=x.device)
    with torch.cuda.device(x.device):
        _packed_silu_mul_kernel[(rows, triton.cdiv(hidden, 1024))](
            out,
            x,
            rows,
            row_stride,
            D=hidden,
            BLOCK=1024,
        )
    return out

View on GitHub (pinned to 0132848349)

Solutions

1. Materialize a clean layout: x = x.contiguous() before the call
2. Ensure x.shape[-1] == 2 * hidden (even) and the tensor is non-empty
3. Guard with the can-use predicate and fall back to eager F.silu(hidden) * gate

Example fix

// before
y = fused_packed_silu_mul_bitexact(x_view)  # strided slice
// after
x_packed = x_view.contiguous()
y = fused_packed_silu_mul_bitexact(x_packed)

ok = (x.is_cuda and x.stride(-1) == 1 and x.stride(-2) >= x.shape[-1]
      and x.stride(0) == x.shape[1] * x.stride(1)
      and x.shape[-1] % 2 == 0 and x.numel() > 0)
if not ok:
    x = x.contiguous()

def usable_packed_silu_mul(x) -> bool:
    return (x.is_cuda and x.stride(-1) == 1
            and x.stride(-2) >= x.shape[-1]
            and x.stride(0) == x.shape[1] * x.stride(1)
            and x.shape[-1] % 2 == 0 and x.numel() > 0)

try:
    y = fused_packed_silu_mul_bitexact(x)
except RuntimeError:
    h, g = x[..., :x.shape[-1]//2].float(), x[..., x.shape[-1]//2:].float()
    y = (h * torch.nn.functional.silu(g)).to(x.dtype)

When it happens

Trigger: Passing a non-contiguous-last-dim view, an odd hidden*2 size, an empty tensor, or a tensor whose leading-dim strides don't satisfy the uniform grouping requirement (e.g. after arbitrary transposes/padding).

Common situations: Feeding SwiGLU projections sliced from a fused MLP buffer with padded strides, or batched 3-D activations whose outer stride was broken by expand/repeat.

Related errors

• kv-canary: {name} must be contiguous
• {fn_name}: dst entry dims (dims {entry_start_dim}..{dst.ndim
• fuse_swiglu_interleaved set on an incompatible fused_moe cal
• The pointers must be multiple of 16 bytes.
• {self._op_label()}: no triton backend

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