{"record":{"id":"5c648c902b01dfba","repo":"xai-org/x-algorithm","slug":"unexpected-swizzle-shift-want-s-3-for-m-4","errorCode":null,"errorMessage":"Unexpected swizzle shift – want S==3 for M==4","messagePattern":"Unexpected swizzle shift – want S==3 for M==4","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":214,"sourceCode":"        Major.K if op.a_major_mode == cute.nvgpu.tcgen05.mma.OperandMajorMode.K else Major.MN,\n        Major.K if op.b_major_mode == cute.nvgpu.tcgen05.mma.OperandMajorMode.K else Major.MN,\n    )\n\n\nclass LayoutType(IntEnum):\n    SWIZZLE_NONE = 0\n    SWIZZLE_128B_BASE32B = 1\n    SWIZZLE_128B = 2\n    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","sourceCodeStart":196,"sourceCodeEnd":232,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L196-L232","documentation":"_layout_type converts a cute.Swizzle into the SMEM descriptor LayoutType. For swizzles with num_base M==4 (the common 128B swizzle family), hardware requires num_shift S==3; any other shift raises ValueError.","triggerScenarios":"Calling make_smem_desc_base with a Swizzle whose num_base is 4 but num_shift is not 3, e.g. Swizzle<3,4,2> or Swizzle<3,4,4>, typically from a hand-built smem layout.","commonSituations":"Constructing custom shared-memory layouts for tcgen05 MMA operands and copying swizzle parameters from SM90 layouts, or using cute.make_swizzle with wrong base/shift arguments.","solutions":["Use a canonical swizzle for M==4: Swizzle<B,4,3> with B in {0,1,2,3} (none/32B/64B/128B)","Build the layout via cute algorithms (e.g. tile_to_shape with a standard swizzle atom) instead of hand-specifying swizzle params","Double-check the Swizzle constructor argument order (bits, base, shift) — a swapped base/shift is the usual cause"],"exampleFix":"# before\nsw = cute.Swizzle(3, 4, 2)  # bad shift\n# after\nsw = cute.Swizzle(3, 4, 3)  # SWIZZLE_128B","handlingStrategy":"validation","validationCode":"B, M, S = sw.num_bits, sw.num_base, sw.num_shift\nif M == 4:\n    assert S == 3, f'Swizzle<,{M},> requires shift 3, got {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":["Copy swizzle atoms from canonical cute definitions instead of constructing Swizzle manually","Remember the constructor order is (num_bits, num_base, num_shift) — swap bugs are common"],"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"}