{"record":{"id":"261983037b946e95","repo":"xai-org/x-algorithm","slug":"not-a-canonical-umma-mn-layout-expected-stride-fa","errorCode":null,"errorMessage":"Not a canonical UMMA_MN Layout: Expected stride failure.","messagePattern":"Not a canonical UMMA_MN Layout: Expected stride failure\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":251,"sourceCode":"    LBO_MODE = 0\n    BASE_OFFSET = 0\n\n    swizzle_atom_mn_size = {\n        LayoutType.SWIZZLE_NONE: 1,\n        LayoutType.SWIZZLE_32B: 2,\n        LayoutType.SWIZZLE_64B: 4,\n        LayoutType.SWIZZLE_128B: 8,\n        LayoutType.SWIZZLE_128B_BASE32B: 8,\n    }[layout_type]\n\n    if major is Major.MN:\n        swizzle_atom_k_size = 4 if layout_type is LayoutType.SWIZZLE_128B_BASE32B else 8\n        canonical_layout = cute.logical_divide(layout, (swizzle_atom_mn_size, swizzle_atom_k_size))\n        if not cute.is_congruent(canonical_layout, ((1, 1), (1, 1))):\n            raise ValueError(\"Not a canonical UMMA_MN Layout: Expected profile failure.\")\n        stride_00 = canonical_layout.stride[0][0]\n        if layout_type is not LayoutType.SWIZZLE_NONE and stride_00 != 1:\n            raise ValueError(\"Not a canonical UMMA_MN Layout: Expected stride failure.\")\n        stride_10 = canonical_layout.stride[1][0]\n        if stride_10 != swizzle_atom_mn_size:\n            raise ValueError(\"Not a canonical UMMA_MN Layout: Expected stride failure.\")\n        stride_01, stride_11 = canonical_layout.stride[0][1], canonical_layout.stride[1][1]\n        if layout_type is LayoutType.SWIZZLE_NONE:\n            stride_byte_offset, leading_byte_offset = stride_01, stride_11\n        else:\n            stride_byte_offset, leading_byte_offset = stride_11, stride_01\n    else:\n        if layout_type == LayoutType.SWIZZLE_128B_BASE32B:\n            raise ValueError(\"SWIZZLE_128B_BASE32B is invalid for Major-K\")\n        if not cute.size(layout.shape[0]) % 8 == 0:\n            raise ValueError(\"Not a canonical UMMA_K Layout: Expected MN-size multiple of 8.\")\n        canonical_layout = cute.logical_divide(layout, (8, 2))\n        if not cute.is_congruent(canonical_layout, ((1, 1), (1, 1))):\n            raise ValueError(\"Not a canonical UMMA_K Layout: Expected profile failure.\")\n        stride_00 = canonical_layout.stride[0][0]\n        if stride_00 != swizzle_atom_mn_size:","sourceCodeStart":233,"sourceCodeEnd":269,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L233-L269","documentation":"After the congruence check passes, make_smem_desc_base verifies that the stride of the innermost MN mode of the swizzle atom is 1 when a swizzle is active. A swizzled MN-major atom must be contiguous in the MN dimension (stride_00 == 1) so the swizzle function can permute the low address bits; anything else cannot be encoded in the descriptor's stride fields.","triggerScenarios":"MN-major layout with a swizzle (SWIZZLE_32B/64B/128B/BASE32B) where the innermost MN stride is not 1 — e.g. a transposed atom where the fastest-varying mode is K, or an atom built with stride (atom_k, 1) instead of (1, atom_k).","commonSituations":"Swapping M/N-major and K-major conventions when porting a kernel; reusing a K-major smem layout while declaring Major.MN; layouts derived from coalesce() or composition that changed the inner stride order.","solutions":["Make the innermost MN stride 1: build the atom as make_layout((atom_mn,atom_k), stride=(1,atom_k)) before tile_to_shape","If the operand really is K-major, pass Major.K instead of Major.MN","If you genuinely want non-contiguous MN, use SWIZZLE_NONE (Swizzle<0,4,3>) which skips this check"],"exampleFix":"# before\natom = cute.make_layout((8,8), stride=(8,1))\n# after\natom = cute.make_layout((8,8), stride=(1,8))","handlingStrategy":"validation","validationCode":"cl = cute.logical_divide(layout, (atom_mn, atom_k))\nif layout_type != LayoutType.SWIZZLE_NONE:\n    assert cl.stride[0][0] == 1, 'inner MN stride must be 1 for swizzled MN-major layouts'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Keep innermost MN stride 1 in swizzled MN-major atoms","Double-check Major enum matches the actual stride order of the tensor"],"tags":["cuda","cutlass","cute-dsl","umma","stride-check","smem-descriptor"],"backgroundTag":"layout-validation-failed","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}