{"record":{"id":"9e49b7758a94943f","repo":"sgl-project/sglang","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":"python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py","lineNumber":259,"sourceCode":"        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:\n            raise ValueError(\"SWIZZLE_128B_BASE32B is invalid for Major-K\")\n        if not cute.size(layout.shape[0]) % 8 == 0:\n            raise ValueError(\n                \"Not a canonical UMMA_K Layout: Expected MN-size multiple of 8.\"\n            )","sourceCodeStart":241,"sourceCodeEnd":277,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py#L241-L277","documentation":"After splitting a Major-MN layout into swizzle atoms, make_smem_desc_base requires the inner MN stride (stride_00) to be exactly 1 whenever the layout is swizzled. A non-unit inner stride means the layout is not contiguous within the atom and cannot be encoded.","triggerScenarios":"A Major.MN swizzled layout where the fastest-varying MN element inside each (atom_mn, atom_k) tile has stride != 1 — e.g. a column-major-ish MN mode or a layout with an interleaved stride pattern.","commonSituations":"Passing a K-major-style layout tagged as Major.MN, or constructing the layout with the mode order swapped.","solutions":["Ensure the MN mode inside each atom is stride-1 contiguous (build with cute.make_layout using compact order).","Double-check the Major argument matches the actual layout construction.","Reuse the layout constructors from the shipped SM100 kernels."],"exampleFix":null,"handlingStrategy":"validation","validationCode":"canon = cute.logical_divide(layout, (atom_mn, atom_k))\nif layout_type is not LayoutType.SWIZZLE_NONE:\n    assert canon.stride[0][0] == 1","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Ensure the MN mode is stride-1 inside each atom for Major.MN.","Verify the Major tag matches layout construction order."],"tags":["cutlass","sm100","layout","stride"],"backgroundTag":"invalid-canonical-layout","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}