{"record":{"id":"7af7026dd9e2c373","repo":"sgl-project/sglang","slug":"the-layout-of-mbias-is-wrong","errorCode":null,"errorMessage":"The layout of mBias is wrong","messagePattern":"The layout of mBias is wrong","errorType":"error_code","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_sm100.py","lineNumber":1301,"sourceCode":"                num_bits_per_copy=universal_copy_bits,\n            )\n            tO_shape_dim_1 = sO_layout.outer.shape[1][0] // async_copy_elems\n            tO_layout = cute.make_ordered_layout(\n                (self.num_epilogue_threads // tO_shape_dim_1, tO_shape_dim_1),\n                order=(1, 0),\n            )\n            # So that we don't have to check if we overshoot kBlockM when we store O\n            assert self.m_block_size % tO_layout.shape[0] == 0\n            vO_layout = cute.make_layout((1, async_copy_elems))\n            gmem_tiled_copy_O = cute.make_tiled_copy_tv(\n                atom_universal_copy, tO_layout, vO_layout\n            )\n\n        if const_expr(mBias is not None):\n            bias_layout_enum = cutlass.utils.LayoutEnum.from_tensor(mBias)\n            self.bias_major_mode = bias_layout_enum.mma_major_mode()\n            if const_expr(self.bias_major_mode != tcgen05.OperandMajorMode.K):\n                raise RuntimeError(\"The layout of mBias is wrong\")\n            # (bias_block_size, n_block_size, bias_stage)\n            sBias_layout = sm100_utils_basic.make_smem_layout_epi(\n                self.bias_dtype,\n                bias_layout_enum,\n                (self.bias_block_size, self.n_block_size),\n                self.bias_stage,\n            )\n            sBias_size = cute.cosize(sBias_layout)\n            # Set after the Q/K/V cta_group_size scaling loop above so bias (non-multicast) isn't double-scaled.\n            self.tma_copy_bytes[\"bias\"] = cute.size_in_bytes(\n                self.bias_dtype, cute.select(sBias_layout, mode=[0, 1])\n            )\n            tma_atom_bias, mBias = cpasync.make_tiled_tma_atom(\n                cpasync.CopyBulkTensorTileG2SOp(),\n                mBias,\n                cute.select(sBias_layout, mode=[0, 1]),\n                (self.bias_block_size, self.n_block_size),\n                1,  # no mcast","sourceCodeStart":1283,"sourceCodeEnd":1319,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/flash_fwd_sm100.py#L1283-L1319","documentation":"The attention bias tensor passed to the SM90/SM100 kernel must be in K-major layout (row/col-major matching tcgen05 OperandMajorMode.K) so TMA can stream it into shared memory for the MMA. A transposed/mixed bias layout is rejected.","triggerScenarios":"Passing an mBias tensor whose memory layout is not K-major (e.g. transposed bias or a non-standard stride order) with attention bias / ALiBi-style score bias enabled.","commonSituations":"Creating bias tensors via .transpose() or .permute() (producing non-contiguous layouts), or materializing bias in (d_k, q) orientation instead of (q, k).","solutions":["Materialize the bias contiguous in K-major (row-major over the K dimension): ensure last-dim stride 1 along the K/head-k axis","Avoid .transpose()/.T on the bias; build it directly in the correct orientation","Call .contiguous() after any transpose"],"exampleFix":"// before\nbias = raw_bias.T  # q-major after transpose\nattn(..., bias=bias)\n// after\nbias = raw_bias.T.contiguous()  # or build K-major directly\nattn(..., bias=bias)","handlingStrategy":"validation","validationCode":"assert bias.stride(-1) == 1 and is_bias_k_major(bias)  # K axis last & contiguous","typeGuard":"def bias_k_major(bias) -> bool:\n    # K-major: last dim (K) has stride 1 and tensor is contiguous\n    return bias.is_contiguous() and bias.stride(-1) == 1","tryCatchPattern":null,"preventionTips":["Construct bias tensors directly in K-major orientation rather than transposing","Call .contiguous() after any transpose of a bias tensor before passing it"],"tags":["flash-attention","tensor-layout","bias","sm100"],"backgroundTag":"tensor-layout-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}