{"record":{"id":"ef721466fe964aa8","repo":"xai-org/x-algorithm","slug":"not-a-canonical-umma-k-layout-expected-profile-fa","errorCode":null,"errorMessage":"Not a canonical UMMA_K Layout: Expected profile failure.","messagePattern":"Not a canonical UMMA_K Layout: Expected profile failure\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":267,"sourceCode":"        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\n    desc |= (int(layout_type) & 0x7) << 61\n\n    return desc & 0xFFFF_FFFF_FFFF_FFFF","sourceCodeStart":249,"sourceCodeEnd":285,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L249-L285","documentation":"K-major counterpart of error 371: after logical_divide(layout,(8,2)), the result must be congruent to ((1,1),(1,1)) — a strict hierarchical layout with an 8x2 atom and uniform strides above it. Any interleaving or non-divisible K extent (K not a multiple of 2) breaks the profile and is rejected before descriptor bit packing.","triggerScenarios":"K-major layout with K extent not divisible by 2, mixed stride hierarchy after dividing by (8,2), or a layout composed with a non-affine value layout so the divide leaves extra modes.","commonSituations":"Odd K sizes (K=1 or K=3 test cases); smem layouts with bank-conflict padding inserted manually that breaks the (8,2) tiling; layouts from cute.composition with swizzle applied before logical_divide.","solutions":["Ensure K is even and the layout is a clean tile_to_shape of an (8,2) K-major atom","Remove manual padding/offsets; let the swizzle handle bank conflicts","If a non-canonical layout is intentional, restructure it as ((8,m),(2,k)) with hierarchical strides before calling"],"exampleFix":"# before\nbad = cute.make_layout((8,3), stride=(1,8))  # K=3 odd\n# after\ngood = cute.make_layout((8,4), stride=(1,8))","handlingStrategy":"validation","validationCode":"assert cute.size(layout.shape[1]) % 2 == 0\nassert cute.is_congruent(cute.logical_divide(layout,(8,2)), ((1,1),(1,1)))","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Keep K even; build layouts by tiling an (8,2) K-major atom","Never insert manual padding into smem layouts destined for UMMA descriptors"],"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"}