{"record":{"id":"5198f3ddcde38828","repo":"xai-org/x-algorithm","slug":"not-a-canonical-umma-mn-layout-expected-profile-f","errorCode":null,"errorMessage":"Not a canonical UMMA_MN Layout: Expected profile failure.","messagePattern":"Not a canonical UMMA_MN Layout: Expected profile failure\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":248,"sourceCode":"    layout_type = _layout_type(swizzle)\n\n    VERSION = 1\n    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))):","sourceCodeStart":230,"sourceCodeEnd":266,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L230-L266","documentation":"When building the UMMA smem descriptor for an MN-major operand, make_smem_desc_base logically divides the layout into the swizzle atom tile ((swizzle_atom_mn_size, 8) — or K-size 4 for 128B_BASE32B) and requires the result to be congruent to ((1,1),(1,1)), i.e. a perfectly hierarchical row-major-within-atom layout. If the divided layout still has interleaved/mixed strides, the layout does not match what the tcgen05 descriptor encoding can express, so it is rejected.","triggerScenarios":"Passing an MN-major cute.Layout whose inner MN/K atom is not a compact rectangular tile of the swizzle atom size — e.g. stride order swapped, a nested layout like ((8,m),(8,k)) with non-uniform strides, or K extent not divisible by the atom K-size (8, or 4 for BASE32B).","commonSituations":"Using an smem tensor tiled from a non-standard atom or with padding in K; mixing a 64B swizzle atom with a layout tiled for 128B; K-dimension size (e.g. K=4 with fp8 and 128B swizzle) not divisible by 8; layouts produced by cute.composition with an arbitrary value layout.","solutions":["Rebuild the operand smem layout as tile_to_shape of the swizzle atom matching your swizzle mode so the atom tile is exactly (atom_mn, atom_k)","Check that the K extent is divisible by swizzle_atom_k_size (8, or 4 for SWIZZLE_128B_BASE32B) and M/N extent divisible by swizzle_atom_mn_size","Verify major: this branch expects Major.MN — if the tensor is K-major pass Major.K","Inspect cute.logical_divide(layout,(atom_mn,atom_k)) shape/stride hierarchy manually to find the non-hierarchical dimension"],"exampleFix":"# before\nsmem_layout = cute.make_layout((M,K), stride=(1,M))  # K-major-ish atom, fails congruence\n# after\natom = cute.tile_to_shape(cute.make_layout((8,8),stride=(1,64)), (M,K))\ndesc = make_smem_desc_base(atom_layout, cute.Swizzle(3,4,3), Major.MN)","handlingStrategy":"validation","validationCode":"atom_mn = {0:1,1:8,2:8,4:4,6:2}[lt]\natom_k = 4 if lt == 1 else 8\ncl = cute.logical_divide(layout, (atom_mn, atom_k))\nassert cute.is_congruent(cl, ((1,1),(1,1))), 'layout will fail UMMA_MN profile check'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Always generate smem layouts via tile_to_shape from the swizzle atom, never compose strides manually","Keep a helper that asserts is_congruent after logical_divide during development"],"tags":["cuda","cutlass","cute-dsl","umma","layout-validation","smem-descriptor"],"backgroundTag":"layout-validation-failed","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}