{"record":{"id":"1c63354bfa355a32","repo":"sgl-project/sglang","slug":"lse-partial-tensor-must-be-float32","errorCode":null,"errorMessage":"LSE partial tensor must be Float32","messagePattern":"LSE partial tensor must be Float32","errorType":"validation","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_combine.py","lineNumber":225,"sourceCode":"        mO_partial: cute.Tensor,\n        mLSE_partial: cute.Tensor,\n        mO: cute.Tensor,\n        mLSE: Optional[cute.Tensor] = None,\n        cu_seqlens: Optional[cute.Tensor] = None,\n        seqused: Optional[cute.Tensor] = None,\n        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(","sourceCodeStart":207,"sourceCodeEnd":243,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_combine.py#L207-L243","documentation":"The FA4 CuTe combine kernel requires the partial log-sum-exp tensor (mLSE_partial) to be float32. LSE values are used to renormalize attention across splits and must be accumulated in fp32 for numerical correctness, so any other dtype is rejected.","triggerScenarios":"Calling the combine __call__ with mLSE_partial of dtype float16/bfloat16 instead of float32.","commonSituations":"Users allocating all attention buffers with the model dtype (bf16/fp16) for uniformity; or casting LSE tensors when moving between devices.","solutions":["Allocate mLSE_partial as torch.float32","Fix whatever produced the partials if it emitted non-fp32 LSE — the upstream kernel contract is fp32 LSE"],"exampleFix":"// before\nmLSE_partial = torch.empty(sp.shape, dtype= torch.bfloat16)\n// after\nmLSE_partial = torch.empty(lse_shape, dtype=torch.float32)","handlingStrategy":"validation","validationCode":"assert mLSE_partial.dtype == torch.float32","typeGuard":"def is_fp32(t) -> bool: return t.dtype == torch.float32","tryCatchPattern":null,"preventionTips":["Remember LSE is always fp32 by contract in flash attention kernels","Add an assert on LSE dtype in test fixtures for attention paths"],"tags":["flash-attention","dtype-mismatch","numerical","validation"],"backgroundTag":"dtype-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}