{"record":{"id":"7d3092ce1987d857","repo":"xai-org/x-algorithm","slug":"only-swizzle-2-5-2-supported-for-128b-base32b","errorCode":null,"errorMessage":"Only Swizzle<2,5,2> supported for 128B_BASE32B","messagePattern":"Only Swizzle<2,5,2> supported for 128B_BASE32B","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":223,"sourceCode":"    SWIZZLE_64B = 4\n    SWIZZLE_32B = 6\n\n\ndef _layout_type(swizzle: cute.Swizzle) -> LayoutType:\n    B, M, S = swizzle.num_bits, swizzle.num_base, swizzle.num_shift\n\n    if M == 4:\n        if S != 3:\n            raise ValueError(\"Unexpected swizzle shift – want S==3 for M==4\")\n        return {\n            0: LayoutType.SWIZZLE_NONE,\n            1: LayoutType.SWIZZLE_32B,\n            2: LayoutType.SWIZZLE_64B,\n            3: LayoutType.SWIZZLE_128B,\n        }[B]\n    if M == 5:\n        if (B, S) != (2, 2):\n            raise ValueError(\"Only Swizzle<2,5,2> supported for 128B_BASE32B\")\n        return LayoutType.SWIZZLE_128B_BASE32B\n\n    raise ValueError(\"Unsupported swizzle triple for UMMA smem descriptor\")\n\n\ndef make_smem_desc_base(layout: cute.Layout, swizzle: cute.Swizzle, major: Major) -> int:\n    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,","sourceCodeStart":205,"sourceCodeEnd":241,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L205-L241","documentation":"_layout_type maps Swizzle<2,5,2> to the special SWIZZLE_128B_BASE32B SMEM layout mode. For swizzles with num_base M==5, only the exact triple (B,S)==(2,2) is supported; anything else raises ValueError.","triggerScenarios":"Calling make_smem_desc_base with a Swizzle whose num_base is 5 but with (num_bits, num_shift) != (2, 2), e.g. Swizzle<3,5,2> or Swizzle<2,5,3>.","commonSituations":"Enabling the 128B_BASE32B layout for fp8/MX-format operands and getting the swizzle atom parameters wrong, or adapting an SM90 swizzle where base-5 atoms have other legal bit widths.","solutions":["Use exactly Swizzle<2,5,2> when you need SWIZZLE_128B_BASE32B","If you need a different swizzle width, switch to an M==4 atom (Swizzle<B,4,3>) which supports 32/64/128B modes","Source the layout from a standard swizzle atom copy (e.g. copy(swizzle_atom)) rather than constructing Swizzle manually"],"exampleFix":"# before\nsw = cute.Swizzle(3, 5, 2)\n# after\nsw = cute.Swizzle(2, 5, 2)  # only legal M==5 triple","handlingStrategy":"validation","validationCode":"B, M, S = sw.num_bits, sw.num_base, sw.num_shift\nif M == 5:\n    assert (B, S) == (2, 2), f'M==5 requires Swizzle<2,5,2>, got Swizzle<{B},{M},{S}>'","typeGuard":"def is_valid_umma_swizzle(sw) -> bool:\n    if sw.num_base == 4:\n        return sw.num_shift == 3 and 0 <= sw.num_bits <= 3\n    if sw.num_base == 5:\n        return (sw.num_bits, sw.num_shift) == (2, 2)\n    return False","tryCatchPattern":null,"preventionTips":["Use Swizzle<2,5,2> whenever enabling 128B_BASE32B mode (fp8/MX operands)","Validate the swizzle triple when building smem descriptors in tests for every layout you configure"],"tags":["cutlass","tcgen05","swizzle","smem-layout","validation"],"backgroundTag":"invalid-swizzle-parameters","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}