{"record":{"id":"d094a34cd7bf794b","repo":"sgl-project/sglang","slug":"o-partial-tensor-must-have-4-or-5-dimensions-num","errorCode":null,"errorMessage":"O partial tensor must have 4 or 5 dimensions: (num_splits, batch, seqlen, nheads, headdim) or (num_splits, total_q, nheads, headdim)","messagePattern":"O partial tensor must have 4 or 5 dimensions: \\(num_splits, batch, seqlen, nheads, headdim\\) or \\(num_splits, total_q, nheads, headdim\\)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_combine.py","lineNumber":231,"sourceCode":"        num_splits_dynamic_ptr: Optional[cute.Tensor] = None,\n        varlen_batch_idx: Optional[cute.Tensor] = None,\n        semaphore_to_reset: Optional[cute.Tensor] = None,\n        # Always keep stream as the last parameter (EnvStream: obtained implicitly via TVM FFI).\n        stream: cuda.CUstream = None,\n    ):\n        # Type checking\n        if const_expr(not (mO_partial.element_type == self.dtype_partial)):\n            raise TypeError(\"O partial tensor must match dtype_partial\")\n        if const_expr(not (mO.element_type == self.dtype)):\n            raise TypeError(\"O tensor must match dtype\")\n        if const_expr(mLSE_partial.element_type not in [Float32]):\n            raise TypeError(\"LSE partial tensor must be Float32\")\n        if const_expr(mLSE is not None and mLSE.element_type not in [Float32]):\n            raise TypeError(\"LSE tensor must be Float32\")\n\n        # Shape validation - input tensors are in user format, need to be converted to kernel format\n        if const_expr(len(mO_partial.shape) not in [4, 5]):\n            raise ValueError(\n                \"O partial tensor must have 4 or 5 dimensions: (num_splits, batch, seqlen, nheads, headdim) or (num_splits, total_q, nheads, headdim)\"\n            )\n        if const_expr(len(mLSE_partial.shape) not in [3, 4]):\n            raise ValueError(\n                \"LSE partial tensor must have 3 or 4 dimensions: (num_splits, batch, seqlen, nheads) or (num_splits, total_q, nheads)\"\n            )\n        if const_expr(len(mO.shape) not in [3, 4]):\n            raise ValueError(\n                \"O tensor must have 3 or 4 dimensions: (batch, seqlen, nheads, headdim) or (total_q, nheads, headdim)\"\n            )\n        if const_expr(mLSE is not None and len(mLSE.shape) not in [2, 3]):\n            raise ValueError(\n                \"LSE tensor must have 2 or 3 dimensions: (batch, seqlen, nheads) or (total_q, nheads)\"\n            )\n\n        mO_partial, mO = [assume_tensor_aligned(t) for t in (mO_partial, mO)]\n        # (num_splits, b, seqlen, h, d) -> (seqlen, d, num_splits, h, b)\n        # or (num_splits, total_q, h, d) -> (total_q, d, num_splits, h)","sourceCodeStart":213,"sourceCodeEnd":249,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_combine.py#L213-L249","documentation":"The combine kernel accepts O partial tensors only in 4D (num_splits, batch, seqlen, nheads, headdim) or 5D varlen forms; other ranks cannot be mapped to the kernel's memory layout.","triggerScenarios":"Passing mO_partial with a rank other than 4 or 5, e.g. a 3D (batch, seqlen, dim) tensor without the num_splits leading axis.","commonSituations":"Feeding raw attention output (3D) instead of per-split partials; dropping the num_splits dimension via squeeze/reshape before combining.","solutions":["Ensure mO_partial keeps the leading num_splits dimension (4D batched or 5D varlen)","If you squeezed it, reshape back: o_partial.unsqueeze(0) when num_splits==1"],"exampleFix":"// before\ncombine(o_3d, lse_partial, mO)  # o_3d: (b, s, h, d)\n// after\ncombine(o_4d, lse_partial, mO)  # o_4d: (num_splits, b, s, h, d)","handlingStrategy":"validation","validationCode":"assert mO_partial.dim() in (4, 5), mO_partial.shape","typeGuard":"def valid_o_partial(t) -> bool: return t.dim() in (4, 5)","tryCatchPattern":null,"preventionTips":["Keep the leading num_splits axis on partial tensors","Log tensor shapes right before split-K combine in debug builds"],"tags":["flash-attention","shape-mismatch","rank-error","validation"],"backgroundTag":"tensor-shape-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}