{"record":{"id":"f9a9f8f5747363ca","repo":"xai-org/x-algorithm","slug":"the-layout-of-mdv-is-wrong","errorCode":null,"errorMessage":"The layout of mdV is wrong","messagePattern":"The layout of mdV is wrong","errorType":"validation","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/flash_bwd_sm100.py","lineNumber":540,"sourceCode":"        self._setup_smem_layout()\n\n        self.cluster_shape_mnk = (*self.cluster_shape_mn, 1)\n        self.cluster_layout_vmnk = cute.tiled_divide(\n            cute.make_layout(self.cluster_shape_mnk),\n            (self.tiled_mma_S.thr_id.shape,),\n        )\n        self.num_mcast_ctas_b = cute.size(self.cluster_layout_vmnk.shape[1])\n        self.is_q_do_mcast = self.num_mcast_ctas_b > 1\n\n        if const_expr(not self.dKV_postprocess):\n            self.mdK_layout_enum = LayoutEnum.from_tensor(mdK)\n            self.mdV_layout_enum = LayoutEnum.from_tensor(mdV)\n            dK_major_mode = self.mdK_layout_enum.mma_major_mode()\n            dV_major_mode = self.mdV_layout_enum.mma_major_mode()\n            if const_expr(dK_major_mode != tcgen05.OperandMajorMode.K):\n                raise RuntimeError(\"The layout of mdK is wrong\")\n            if const_expr(dV_major_mode != tcgen05.OperandMajorMode.K):\n                raise RuntimeError(\"The layout of mdV is wrong\")\n\n        if const_expr(self.use_tma_store and not self.dKV_postprocess):\n            tma_copy_op_dKV = cpasync.CopyBulkTensorTileS2GOp()\n            tma_atom_dK, mdK_tma_tensor = cpasync.make_tiled_tma_atom(\n                tma_copy_op_dKV,\n                mdK,\n                cute.select(self.sdK_layout, mode=[0, 1]),\n                self.sdK_epi_tile,\n                1,\n            )\n            tma_atom_dV, mdV_tma_tensor = cpasync.make_tiled_tma_atom(\n                tma_copy_op_dKV,\n                mdV,\n                cute.select(self.sdV_layout, mode=[0, 1]),\n                self.sdV_epi_tile,\n                1,\n            )\n        else:","sourceCodeStart":522,"sourceCodeEnd":558,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/flash_bwd_sm100.py#L522-L558","documentation":"In the SM100 flash-attention backward kernel, mdV (the dV accumulator tensor) is expected to be laid out in K-major (column-major) mode as required by tcgen05 MMA operands. The kernel derives the layout via LayoutEnum.from_tensor(mdV) and checks its mma_major_mode(); if it is not OperandMajorMode.K it aborts because the MMA instruction descriptor would be invalid.","triggerScenarios":"Calling the flash backward kernel __call__ with an mdV tensor whose memory layout is M-major/row-major (e.g. created transposed or with a swapped stride order) instead of K-major, so mma_major_mode() returns MN mode.","commonSituations":"Passing dV output tensors allocated with wrong stride order, reusing tensors from the forward pass without transposing, or changes in how the host-side wrapper constructs mdV after a refactor of the kernel signature.","solutions":["Check how mdV is allocated/passed on the host side and ensure its contiguous dimension is the K (head-dim/inner) axis so the layout enum resolves to K-major","Transpose or re-layout mdV (cute.select or .T equivalent) before invoking the kernel, mirroring how mdK is prepared","Verify mdK and mdV are constructed identically; the mdK check above passing while mdV fails points at asymmetry in their creation","If tensors come from PyTorch, ensure the corresponding dV tensor is made contiguous in the same dim as dK"],"exampleFix":"// before\nmdV = cute.make_tensor(dV_ptr, mD_layout)  # M-major\n// after\nmdV = cute.make_tensor(dV_ptr, cute.select(mD_layout, mode=[1, 0, 2, 3]))  # K-major","handlingStrategy":"validation","validationCode":"from phoenix.xrex.cutedsl.ranker_fa4.flash_bwd_sm100 import LayoutEnum\nimport cutlass.tcgen05 as tcgen05\nassert LayoutEnum.from_tensor(mdV).mma_major_mode() == tcgen05.OperandMajorMode.K, 'mdV must be K-major'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Construct dK and dV tensors with identical layouts so the K/V paths stay symmetric","Wrap kernel launches in a debug helper that asserts both mdK and mdV major modes before calling __call__"],"tags":["cuda","flash-attention","layout","tcgen05","sm100"],"backgroundTag":"tensor-layout-mismatch","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}