{"record":{"id":"ce8ac7e6fe7d5d32","repo":"sgl-project/sglang","slug":"the-layout-of-v-is-not-supported","errorCode":null,"errorMessage":"The layout of v is not supported","messagePattern":"The layout of v 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":431,"sourceCode":"            )\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,\n            self.qk_acc_dtype,\n            cta_group,","sourceCodeStart":413,"sourceCodeEnd":449,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/sm100_hd256_2cta_fmha_forward.py#L413-L449","documentation":"The kernel requires V in MN-major (column-major) layout, the opposite of Q/K, per the second operand of the PV MMA. LayoutEnum.from_tensor(v).mma_major_mode() returned K-major, so V was passed row-major.","triggerScenarios":"Passing a fully contiguous/row-major v tensor where the kernel expects MN-major layout for V.","commonSituations":"Calling .contiguous() uniformly on q/k/v 'for safety'; V coming from a cache that stores V in the same layout as K.","solutions":["Provide v in MN-major layout as produced by the intended cache/prepare path","Review how the kernel's reference caller prepares v (e.g. v.transpose or cache layout) and mirror it","Do not blanket-contiguous() v before this kernel"],"exampleFix":"# before\nout = fmha(q, k, v.contiguous())\n# after\nout = fmha(q, k, v)  # keep v in the kernel-expected MN-major layout","handlingStrategy":"type-guard","validationCode":"assert v.stride(-2) == 1 or layout_matches_kernel(v), 'v must be MN-major'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Do not blanket-contiguous() v; follow kernel's layout docs","Keep a reference invocation in tests to catch layout drift"],"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"}