{"record":{"id":"b5822535a9ac9285","repo":"sgl-project/sglang","slug":"unsupported-input-for-sana-fused-bias-glu","errorCode":null,"errorMessage":"unsupported input for Sana fused bias-GLU","messagePattern":"unsupported input for Sana fused bias-GLU","errorType":"exception","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/activation/sana_conv_post_triton.py","lineNumber":102,"sourceCode":"    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()\n    )\n\n\ndef fused_bias_glu(x: torch.Tensor, bias: torch.Tensor) -> torch.Tensor:\n    if not can_use_fused_bias_glu(x, bias):\n        raise RuntimeError(\"unsupported input for Sana fused bias-GLU\")\n    batch, double_channels, height, width = x.shape\n    channels = double_channels // 2\n    out = torch.empty(\n        (batch, channels, height, width),\n        dtype=x.dtype,\n        device=x.device,\n        memory_format=torch.channels_last,\n    )\n    with torch.cuda.device(x.device):\n        _bias_glu_kernel[(triton.cdiv(out.numel(), 1024),)](\n            out, x, bias, out.numel(), channels=channels\n        )\n    return out\n\n\n__all__ = [\n    \"can_use_fused_bias_glu\",\n    \"can_use_fused_bias_silu\",","sourceCodeStart":84,"sourceCodeEnd":120,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/activation/sana_conv_post_triton.py#L84-L120","documentation":"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').","triggerScenarios":"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).","commonSituations":"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.","solutions":["Ensure channels are even and tensor is channels_last bf16: x = x.contiguous(memory_format=torch.channels_last).bfloat16()","Verify bias.shape == (x.shape[1],) and bias is on x.device with bf16 dtype","Guard with can_use_fused_bias_glu(x, bias) and use an eager fallback"],"exampleFix":"// before\ny = fused_bias_glu(x, bias)\n// after\nif can_use_fused_bias_glu(x, bias):\n    y = fused_bias_glu(x, bias)\nelse:\n    h, g = x[:, :C].float(), x[:, C:].float()\n    y = ((h + bias[:C, None, None]).float() * torch.sigmoid(g + bias[C:, None, None]).float()).to(x.dtype)","handlingStrategy":"fallback","validationCode":"from sglang.kernels.ops.diffusion.activation.sana_conv_post_triton import can_use_fused_bias_glu\nif not can_use_fused_bias_glu(x, bias):\n    raise ValueError('fall back to eager GLU')  # or route to eager path","typeGuard":"def usable_bias_glu(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 x.shape[1] % 2 == 0\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_glu(x, bias)\nexcept RuntimeError:\n    C = x.shape[1] // 2\n    y = eager_bias_glu(x, bias, C)","preventionTips":["Assert x.shape[1] % 2 == 0 before the GLU projection","Keep the whole Sana pipeline in channels_last bf16 to satisfy the fused contract"],"tags":["sana","diffusion","glu","channels-last","triton"],"backgroundTag":"unsupported-kernel-input-contract","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}