{"record":{"id":"96d381f266455a5d","repo":"sgl-project/sglang","slug":"the-layout-of-q-is-not-supported","errorCode":null,"errorMessage":"The layout of q is not supported","messagePattern":"The layout of q is not supported","errorType":"exception","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py","lineNumber":427,"sourceCode":"            self.tile_sched_params, grid = compute_grid_clc(\n                (s_q, o.shape[1], o.shape[2]) if cum_seqlen_q is not None else o.shape,\n                self.cta_tiler,\n                (*self.cluster_shape_mn, 1),\n            )\n        else:\n            self.tile_sched_params, grid = compute_grid(\n                (s_q, o.shape[1], o.shape[2]) if cum_seqlen_q is not None else o.shape,\n                self.cta_tiler,\n                self.is_persistent,\n            )\n\n        self.q_major_mode = utils.LayoutEnum.from_tensor(q).mma_major_mode()\n        self.k_major_mode = utils.LayoutEnum.from_tensor(k).mma_major_mode()\n        self.v_major_mode = utils.LayoutEnum.from_tensor(v).mma_major_mode()\n        self.o_layout = utils.LayoutEnum.from_tensor(o)\n\n        if cutlass.const_expr(self.q_major_mode != tcgen05.OperandMajorMode.K):\n            raise RuntimeError(\"The layout of q is not supported\")\n        if cutlass.const_expr(self.k_major_mode != tcgen05.OperandMajorMode.K):\n            raise RuntimeError(\"The layout of k is not supported\")\n        if cutlass.const_expr(self.v_major_mode != tcgen05.OperandMajorMode.MN):\n            raise RuntimeError(\"The layout of v is not supported\")\n\n        # check type consistency\n        if cutlass.const_expr(self.q_dtype != self.k_dtype):\n            raise TypeError(f\"Type mismatch: {self.q_dtype} != {self.k_dtype}\")\n        if cutlass.const_expr(self.q_dtype != self.v_dtype):\n            raise TypeError(f\"Type mismatch: {self.q_dtype} != {self.v_dtype}\")\n        self._setup_attributes()\n\n        cta_group = tcgen05.CtaGroup.TWO\n        # the intermediate tensor p is from tmem & k-major\n        p_source = tcgen05.OperandSource.TMEM\n        p_major_mode = tcgen05.OperandMajorMode.K\n        qk_tiled_mma = sm100_utils.make_trivial_tiled_mma(\n            self.q_dtype,","sourceCodeStart":409,"sourceCodeEnd":445,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py#L409-L445","documentation":"The SM100 hd256 FMHA kernel requires Q in K-major (row-major per head) MMA layout; it inspects the tensor's stride pattern via LayoutEnum.from_tensor(q).mma_major_mode(). A MN-major (column-major) Q is not supported by this tcgen05 kernel configuration.","triggerScenarios":"Passing q whose last dim stride != 1 (transposed/column-major) to the hd256 2cta fmha forward call.","commonSituations":"Upstream code transposes q for a different attention backend; a projection output happens to be non-contiguous; mixing kernels with differing layout requirements.","solutions":["Make q contiguous in the last dim: q = q.contiguous()","Check that qkv projection produces row-major output","Verify no transpose was applied to q before the call"],"exampleFix":"# before\nout = fmha(q.transpose(-1,-2).contiguous(), k, v)  # wrong layout\n# after\nout = fmha(q.contiguous(), k, v)","handlingStrategy":"type-guard","validationCode":"assert q.stride(-1) == 1, 'q must be K-major (row-major)'","typeGuard":"def q_is_k_major(q: torch.Tensor) -> bool:\n    return q.stride(-1) == 1","tryCatchPattern":null,"preventionTips":["Call .contiguous() on q at the wrapper boundary","Never pass transposed q views across attention kernel APIs"],"tags":["cuda","flash-attention","memory-layout","sm100"],"backgroundTag":"unsupported-tensor-layout","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}