{"record":{"id":"11a5450930028154","repo":"xai-org/x-algorithm","slug":"not-a-canonical-umma-k-layout-expected-mn-size-mu","errorCode":null,"errorMessage":"Not a canonical UMMA_K Layout: Expected MN-size multiple of 8.","messagePattern":"Not a canonical UMMA_K Layout: Expected MN-size multiple of 8\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":264,"sourceCode":"        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:\n            raise ValueError(\"Not a canonical UMMA_K Layout: Expected stride failure.\")\n        stride_10 = canonical_layout.stride[1][0]\n        if layout_type is not LayoutType.SWIZZLE_NONE and stride_10 != 1:\n            raise ValueError(\"Not a canonical UMMA_K Layout: Expected stride failure.\")\n        stride_01 = canonical_layout.stride[0][1]\n        stride_byte_offset, leading_byte_offset = stride_01, stride_10\n\n    desc = 0\n    desc |= (leading_byte_offset & 0x3FFF) << 16\n    desc |= (stride_byte_offset & 0x3FFF) << 32\n    desc |= (VERSION & 0x3) << 46\n    desc |= (BASE_OFFSET & 0x7) << 49\n    desc |= (LBO_MODE & 0x1) << 52","sourceCodeStart":246,"sourceCodeEnd":282,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L246-L282","documentation":"For K-major operands, make_smem_desc_base requires the MN extent of the smem layout to be a multiple of 8 so it can logical_divide by an (8,2) atom. A non-multiple means the layout cannot be partitioned into the 8x2 core tiles the UMMA descriptor assumes.","triggerScenarios":"K-major layout whose layout.shape[0] (M/N size) is e.g. 64? fine, but 4, 6, 12, or any value not divisible by 8 — typical with small unit-test tiles or tail tiles.","commonSituations":"Unit tests with tiny shapes like (4,K); partial/epilogue tiles of an M dimension not padded to a multiple of 8; ranker kernels whose head dimension was changed to a non-multiple-of-8 value.","solutions":["Pad or choose M/N extent to a multiple of 8 (the tmem/MMA atom size)","Use an MN-major layout for that operand instead, which is validated against the swizzle atom instead of the 8-multiple rule","Ensure the kernel's tile shape M is at least the MMA atom M (128 on tcgen05) so the smem M is naturally a multiple of 8"],"exampleFix":"# before\ndesc = make_smem_desc_base(cute.make_layout((6,K),stride=(K,1)), sw, Major.K)\n# after\ndesc = make_smem_desc_base(cute.make_layout((8,K),stride=(K,1)), sw, Major.K)","handlingStrategy":"validation","validationCode":"assert cute.size(layout.shape[0]) % 8 == 0, 'K-major UMMA operand M/N must be multiple of 8'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Pad M/N smem extents to multiples of 8","Avoid tiny non-multiple-of-8 test tiles or switch them to MN-major"],"tags":["cuda","cutlass","cute-dsl","umma","shape-validation","smem-descriptor"],"backgroundTag":"shape-not-multiple-of-tile","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}