{"record":{"id":"001bfb6a5ba5dba0","repo":"sgl-project/sglang","slug":"unsupported-input-for-sana-fused-bias-silu","errorCode":null,"errorMessage":"unsupported input for Sana fused bias-SiLU","messagePattern":"unsupported input for Sana fused bias-SiLU","errorType":"exception","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/activation/sana_conv_post_triton.py","lineNumber":78,"sourceCode":"        and x.is_contiguous(memory_format=torch.channels_last)\n    )\n\n\ndef can_use_fused_bias_silu(x: torch.Tensor, bias: torch.Tensor) -> bool:\n    return (\n        _is_channels_last_bf16(x)\n        and bias.is_cuda\n        and bias.dtype is x.dtype\n        and bias.device == x.device\n        and bias.dim() == 1\n        and bias.shape[0] == x.shape[1]\n        and bias.is_contiguous()\n    )\n\n\ndef fused_bias_silu(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:\n    if not can_use_fused_bias_silu(x, bias):\n        raise RuntimeError(\"unsupported input for Sana fused bias-SiLU\")\n    out = torch.empty_like(x, memory_format=torch.preserve_format)\n    with torch.cuda.device(x.device):\n        _bias_silu_kernel[(triton.cdiv(x.numel(), 1024),)](\n            out, x, bias, x.numel(), channels=x.shape[1]\n        )\n    return out\n\n\ndef can_use_fused_bias_glu(x: torch.Tensor, bias: torch.Tensor) -> bool:\n    return (\n        _is_channels_last_bf16(x)\n        and x.shape[1] % 2 == 0\n        and bias.is_cuda\n        and bias.dtype is x.dtype\n        and bias.device == x.device\n        and bias.dim() == 1\n        and bias.shape[0] == x.shape[1]\n        and bias.is_contiguous()","sourceCodeStart":60,"sourceCodeEnd":96,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/activation/sana_conv_post_triton.py#L60-L96","documentation":"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').","triggerScenarios":"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.","commonSituations":"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.","solutions":["Convert layout: x = x.to(memory_format=torch.channels_last, dtype=torch.bfloat16)","Match bias: bias = bias.to(x.dtype, x.device); assert bias.shape[0] == x.shape[1]","Guard with can_use_fused_bias_silu(x, bias) and fall back to F.silu(x + bias[:, None, None]) otherwise"],"exampleFix":"// before\ny = fused_bias_silu(x, bias)  # x is channels_first fp32\n// after\nif can_use_fused_bias_silu(x, bias):\n    y = fused_bias_silu(x, bias)\nelse:\n    y = F.silu(x.float() + bias[:, None, None].float()).to(x.dtype)","handlingStrategy":"fallback","validationCode":"from sglang.kernels.ops.diffusion.activation.sana_conv_post_triton import can_use_fused_bias_silu\nif not can_use_fused_bias_silu(x, bias):\n    x = x.to(memory_format=torch.channels_last, dtype=torch.bfloat16)\n    bias = bias.to(torch.bfloat16, x.device)","typeGuard":"def usable_bias_silu(x, b) -> bool:\n    return (x.is_cuda and x.dtype is torch.bfloat16 and x.dim() == 4\n            and x.numel() > 0 and x.is_contiguous(memory_format=torch.channels_last)\n            and b.is_cuda and b.dtype is x.dtype and b.device == x.device\n            and b.dim() == 1 and b.shape[0] == x.shape[1] and b.is_contiguous())","tryCatchPattern":"try:\n    y = fused_bias_silu(x, bias)\nexcept RuntimeError:\n    y = torch.nn.functional.silu(x.float() + bias[:, None, None].float()).to(x.dtype)","preventionTips":["Run Sana convs in bf16 channels_last end-to-end (model.to(memory_format=torch.channels_last))","Always pair the fused call with can_use_fused_bias_silu and an eager fallback"],"tags":["sana","diffusion","triton","memory-format","channels-last"],"backgroundTag":"unsupported-kernel-input-contract","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}