{"record":{"id":"72fbec32c085d9cd","repo":"sgl-project/sglang","slug":"the-layout-of-k-is-not-supported","errorCode":null,"errorMessage":"The layout of k is not supported","messagePattern":"The layout of k 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":429,"sourceCode":"                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,\n            self.q_major_mode,\n            self.k_major_mode,","sourceCodeStart":411,"sourceCodeEnd":447,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py#L411-L447","documentation":"The kernel requires K in K-major layout for the tcgen05 MMA. LayoutEnum.from_tensor(k).mma_major_mode() returned MN-major, meaning K is effectively column-major, which this kernel cannot consume.","triggerScenarios":"Passing a column-major (transposed) k tensor to hd256 2cta fmha forward.","commonSituations":"KV cache stored transposed for another backend (e.g. FA-style [B,H,S,D] vs transposed views); reusing tensors prepared for flashinfer with different layout expectations.","solutions":["Pass k contiguous with unit stride in the head dim: k.contiguous()","Verify the KV cache pool layout matches the kernel expectation (K-major)","Fix upstream reshape/transpose of the K cache"],"exampleFix":"# before\nout = fmha(q, k_cache_view.transpose(-1,-2), v)\n# after\nout = fmha(q, k_cache_view.contiguous(), v)","handlingStrategy":"type-guard","validationCode":"assert k.stride(-1) == 1, 'k must be K-major'","typeGuard":"def k_is_k_major(k: torch.Tensor) -> bool:\n    return k.stride(-1) == 1","tryCatchPattern":null,"preventionTips":["Document the KV cache layout contract per backend","Assert strides in debug builds of the attention wrapper"],"tags":["cuda","flash-attention","memory-layout","sm100","kv-cache"],"backgroundTag":"unsupported-tensor-layout","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}