{"record":{"id":"278103e8d621ae2b","repo":"sgl-project/sglang","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":"python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py","lineNumber":254,"sourceCode":"    LBO_MODE = 0  # bit  52\n    BASE_OFFSET = 0  # bits 49–51   (CUTLASS always 0)\n\n    # ---------------------------------------------------------- strides  (units: uint128_t = 16 B)\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(\n            layout, (swizzle_atom_mn_size, swizzle_atom_k_size)\n        )\n        if not cute.is_congruent(canonical_layout, ((1, 1), (1, 1))):\n            raise ValueError(\n                \"Not a canonical UMMA_MN Layout: Expected profile failure.\"\n            )\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 = (\n            canonical_layout.stride[0][1],\n            canonical_layout.stride[1][1],\n        )\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:","sourceCodeStart":236,"sourceCodeEnd":272,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py#L236-L272","documentation":"make_smem_desc_base validates that a Major-MN shared-memory layout is congruent to the canonical swizzle-atom tiling ((atom_mn, atom_k) split into ((1,1),(1,1))). If cute.is_congruent fails, the layout does not have the profile the descriptor encoder assumes, so it refuses rather than encoding garbage.","triggerScenarios":"Passing a Major.MN layout whose shape/stride structure is not a clean product of (swizzle_atom_mn_size, swizzle_atom_k_size) tiles — e.g. with remainders, nested hierarchies, or stride holes — via gemm_ptx* or smem_desc_base_from_tensor.","commonSituations":"Building smem tile layouts from tensor views with padding or non-uniform strides; composing layouts with cute.tile_to_shape incorrectly so the K or MN remainder leaks into the profile.","solutions":["Construct the layout with cute.make_layout / cute.tile_to_shape so it is an exact product of the swizzle atom (atom_mn x 8, or 4 for 128B_BASE32B).","Compare your layout against the ones built inside the shipped gemm_ptx* helpers and match their construction path.","Falling back to smem_desc_base_from_tensor on a tensor whose cute view already has a canonical layout avoids hand-rolled profiles."],"exampleFix":null,"handlingStrategy":"validation","validationCode":"import cute\n\natom_k = 4 if is_base32b else 8\ncanon = cute.logical_divide(layout, (atom_mn, atom_k))\nassert cute.is_congruent(canon, ((1, 1), (1, 1)))","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Build layouts only via make_layout/tile_to_shape with exact-multiple shapes.","Add congruence unit tests for every layout passed to make_smem_desc_base."],"tags":["cutlass","sm100","layout","shared-memory"],"backgroundTag":"invalid-canonical-layout","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}