{"record":{"id":"da23af0525c0ebe2","repo":"sgl-project/sglang","slug":"this-layer-doesn-t-support-feature-dim-64kb","errorCode":null,"errorMessage":"This layer doesn't support feature dim >= 64KB.","messagePattern":"This layer doesn't support feature dim >= 64KB\\.","errorType":"validation","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/fla/l2norm.py","lineNumber":90,"sourceCode":"\ndef l2norm_fwd(\n    x: torch.Tensor, eps: float = 1e-6, output_dtype: Optional[torch.dtype] = None\n):\n    x_shape_og = x.shape\n    x = x.view(-1, x.shape[-1])\n    # allocate output\n    if output_dtype is None:\n        y = torch.empty_like(x)\n    else:\n        y = torch.empty_like(x, dtype=output_dtype)\n    assert y.stride(-1) == 1\n    T, D = x.shape[0], x.shape[-1]\n    # rstd = torch.empty((T,), dtype=torch.float32, device=x.device)\n    # Less than 64KB per feature: enqueue fused kernel\n    MAX_FUSED_SIZE = 65536 // x.element_size()\n    BD = min(MAX_FUSED_SIZE, triton.next_power_of_2(D))\n    if D > BD:\n        raise RuntimeError(\"This layer doesn't support feature dim >= 64KB.\")\n\n    if D <= 512:\n\n        def grid(meta):\n            return (triton.cdiv(T, meta[\"BT\"]),)\n\n        l2norm_fwd_kernel[grid](\n            x,\n            y,\n            eps,\n            T=T,\n            D=D,\n            BD=BD,\n            BT=16,\n            num_warps=8,\n            num_stages=3,\n        )\n    else:","sourceCodeStart":72,"sourceCodeEnd":108,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/fla/l2norm.py#L72-L108","documentation":"l2norm_fwd uses a Triton fused kernel that loads an entire feature row into one block; the maximum fused size is 65536 bytes divided by element size. If the feature dimension D exceeds that (e.g. >16384 for fp32 or >32768 for bf16), no kernel path exists and it raises RuntimeError.","triggerScenarios":"Calling l2norm_fwd (via chunk_kda, forward, or cutedsl/flashinfer wrappers) with x.shape[-1] * element_size > 65536 bytes.","commonSituations":"Using an unusually large head_dim or hidden dim (e.g. D=32768 in fp32, D=65536+ in bf16) in KDA/attention normalization layers.","solutions":["Reduce the feature dimension to below 64KB per row (e.g. <16384 for fp32)","Split the tensor along the feature dim, l2norm each chunk, and concatenate","Use a smaller dtype (bf16/fp16 halves the byte limit)"],"exampleFix":"// before\ny = l2norm_fwd(x)  # x: [T, 65536] fp32 -> raises\n// after\nx = x.to(torch.bfloat16)  # or split feature dim\ny = l2norm_fwd(x)","handlingStrategy":"validation","validationCode":"assert x.shape[-1] * x.element_size() <= 65536, f'feature dim {x.shape[-1]} exceeds 64KB fused-kernel limit'","typeGuard":null,"tryCatchPattern":"try:\\n    y = l2norm_fwd(x)\\nexcept RuntimeError:\\n    x = x.to(torch.bfloat16); y = l2norm_fwd(x)","preventionTips":["Check D * element_size <= 65536 before calling","Prefer bf16 activations for large head dims"],"tags":["triton","l2norm","feature-dim","kda","attention"],"backgroundTag":"kernel-feature-dim-limit-exceeded","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}