{"record":{"id":"d11d5e2fbe39aaad","repo":"sgl-project/sglang","slug":"unsupported-input-for-packed-fused-silu-mul","errorCode":null,"errorMessage":"unsupported input for packed fused SiLU-mul","messagePattern":"unsupported input for packed fused SiLU-mul","errorType":"exception","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/activation/silu_mul_bitexact.py","lineNumber":113,"sourceCode":"            numel,\n            BLOCK=1024,\n        )\n    return out\n\n\ndef fused_packed_silu_mul_bitexact(x: torch.Tensor) -> torch.Tensor:\n    \"\"\"Bit-exact SwiGLU over a contiguous packed ``[..., 2 * D]`` input.\"\"\"\n    if not (\n        x.is_cuda\n        and x.dtype is torch.bfloat16\n        and x.dim() == 3\n        and x.stride(-1) == 1\n        and x.stride(-2) >= x.shape[-1]\n        and x.stride(0) == x.shape[1] * x.stride(1)\n        and x.shape[-1] % 2 == 0\n        and x.numel() > 0\n    ):\n        raise RuntimeError(\"unsupported input for packed fused SiLU-mul\")\n    hidden = x.shape[-1] // 2\n    rows = x.numel() // x.shape[-1]\n    row_stride = x.stride(-2)\n    out = torch.empty((*x.shape[:-1], hidden), dtype=x.dtype, device=x.device)\n    with torch.cuda.device(x.device):\n        _packed_silu_mul_kernel[(rows, triton.cdiv(hidden, 1024))](\n            out,\n            x,\n            rows,\n            row_stride,\n            D=hidden,\n            BLOCK=1024,\n        )\n    return out\n","sourceCodeStart":95,"sourceCodeEnd":128,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/activation/silu_mul_bitexact.py#L95-L128","documentation":"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').","triggerScenarios":"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).","commonSituations":"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.","solutions":["Materialize a clean layout: x = x.contiguous() before the call","Ensure x.shape[-1] == 2 * hidden (even) and the tensor is non-empty","Guard with the can-use predicate and fall back to eager F.silu(hidden) * gate"],"exampleFix":"// before\ny = fused_packed_silu_mul_bitexact(x_view)  # strided slice\n// after\nx_packed = x_view.contiguous()\ny = fused_packed_silu_mul_bitexact(x_packed)","handlingStrategy":"fallback","validationCode":"ok = (x.is_cuda and x.stride(-1) == 1 and x.stride(-2) >= x.shape[-1]\n      and x.stride(0) == x.shape[1] * x.stride(1)\n      and x.shape[-1] % 2 == 0 and x.numel() > 0)\nif not ok:\n    x = x.contiguous()","typeGuard":"def usable_packed_silu_mul(x) -> bool:\n    return (x.is_cuda and x.stride(-1) == 1\n            and x.stride(-2) >= x.shape[-1]\n            and x.stride(0) == x.shape[1] * x.stride(1)\n            and x.shape[-1] % 2 == 0 and x.numel() > 0)","tryCatchPattern":"try:\n    y = fused_packed_silu_mul_bitexact(x)\nexcept RuntimeError:\n    h, g = x[..., :x.shape[-1]//2].float(), x[..., x.shape[-1]//2:].float()\n    y = (h * torch.nn.functional.silu(g)).to(x.dtype)","preventionTips":["Call .contiguous() on SwiGLU projections sliced from padded buffers","Keep hidden*2 as the trailing dim with stride 1"],"tags":["swiglu","silu","triton","strides","bit-exact"],"backgroundTag":"unsupported-kernel-input-contract","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}