{"record":{"id":"657b745795c750ef","repo":"sgl-project/sglang","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":"python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py","lineNumber":280,"sourceCode":"            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            )\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    # ------------------------------------------------------------------ pack\n    desc = 0\n    # leading_byte_offset_  [16:30)\n    desc |= (leading_byte_offset & 0x3FFF) << 16\n    # stride_byte_offset_   [32:46)\n    desc |= (stride_byte_offset & 0x3FFF) << 32\n    # version_             [46:48)\n    desc |= (VERSION & 0x3) << 46\n    # base_offset_         [49:52)","sourceCodeStart":262,"sourceCodeEnd":298,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py#L262-L298","documentation":"For Major-K layouts, make_smem_desc_base splits the layout into (8, 2) atoms and requires the result be congruent to ((1,1),(1,1)) — a clean two-level hierarchy. If cute.is_congruent fails, the layout profile is not canonical and cannot be encoded into the smem descriptor.","triggerScenarios":"A Major.K layout with remainders in either mode, nested sub-modes, or non-uniform strides passed to gemm_ptx* / smem_desc_base_from_tensor.","commonSituations":"Building K-major smem tiles with cute.composition of mismatched atoms; layouts with a leftover mode from tile_to_shape when shape isn't a multiple of the atom.","solutions":["Ensure both layout extents are exact multiples of the (8, 2) atom split; pad if needed.","Rebuild the layout with cute.make_layout / cute.tile_to_shape so no remainder modes survive.","Test cute.is_congruent(cute.logical_divide(layout, (8, 2)), ((1,1),(1,1))) in a unit test before kernel launch."],"exampleFix":null,"handlingStrategy":"validation","validationCode":"canon = cute.logical_divide(layout, (8, 2))\nassert cute.is_congruent(canon, ((1, 1), (1, 1)))","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Make K-major layout extents exact multiples of the (8,2) atom.","Build layouts from canonical atoms; never compose mismatched hierarchies."],"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"}