{"record":{"id":"736bae6eb9541b3c","repo":"xai-org/x-algorithm","slug":"unsupported-swizzle-triple-for-umma-smem-descripto","errorCode":null,"errorMessage":"Unsupported swizzle triple for UMMA smem descriptor","messagePattern":"Unsupported swizzle triple for UMMA smem descriptor","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":226,"sourceCode":"\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,\n    }[layout_type]\n\n    if major is Major.MN:","sourceCodeStart":208,"sourceCodeEnd":244,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L208-L244","documentation":"_layout_type maps a cute.Swizzle<B,M,S> triple to one of the UMMA shared-memory descriptor layout modes (SWIZZLE_NONE/32B/64B/128B/128B_BASE32B). Only M==4 with S==3 (canonical CUDA swizzles) or the exact triple Swizzle<2,5,2> are recognized; any other combination has no hardware encoding in this builder, so it rejects it. The error surfaces from make_smem_desc_base, which is the entry point used by all gemm_ptx kernels and smem_desc_base_from_tensor.","triggerScenarios":"Calling make_smem_desc_base / smem_desc_base_from_tensor / gemm_ptx* with a Swizzle whose num_base (M) is neither 4 nor 5, e.g. Swizzle<3,3,3>; or M==4 with B>3 or S!=3 variants that escape the earlier check; or M==5 with (B,S)!=(2,2) (covered by the adjacent BASE32B error before reaching this line).","commonSituations":"Hand-rolled smem atom layouts copied from a different CUTLASS/CuTe DSL tutorial with a non-standard swizzle; porting Blackwell tcgen05 kernels that use sub-64B or 32B-base swizzle variants not supported here; upgrading the CUTLASS Python DSL where swizzle atom conventions changed.","solutions":["Use a canonical swizzle: Swizzle<0,4,3> (none), <1,4,3> (32B), <2,4,3> (64B), <3,4,3> (128B), or Swizzle<2,5,2> for 128B_BASE32B","Build the smem layout from a tile_to_shape over a standard Swizzle atom (e.g. cute.make_layout of tile_to_shape(SmemLayoutAtom, smem_shape)) instead of composing a custom swizzle","If you need an unsupported swizzle mode, extend the LayoutType enum and the swizzle_atom_mn_size table in this file rather than passing the odd swizzle","Print swizzle.num_bits/num_base/num_shift right before the call to confirm which triple you are actually generating"],"exampleFix":"# before\ndesc = make_smem_desc_base(layout, cute.Swizzle(3,3,3), Major.MN)\n# after\ndesc = make_smem_desc_base(layout, cute.Swizzle(3,4,3), Major.MN)","handlingStrategy":"validation","validationCode":"SUPPORTED = {(0,4,3),(1,4,3),(2,4,3),(3,4,3),(2,5,2)}\ndef is_supported_swizzle(sw):\n    return (sw.num_bits, sw.num_base, sw.num_shift) in SUPPORTED","typeGuard":"def is_supported_swizzle(sw: cute.Swizzle) -> bool:\n    return (sw.num_bits, sw.num_base, sw.num_shift) in {(0,4,3),(1,4,3),(2,4,3),(3,4,3),(2,5,2)}","tryCatchPattern":"try:\n    desc = make_smem_desc_base(layout, sw, major)\nexcept ValueError as e:\n    if 'Unsupported swizzle' in str(e):\n        sw = cute.Swizzle(3,4,3)  # fall back to 128B canonical\n        desc = make_smem_desc_base(layout, sw, major)\n    else:\n        raise","preventionTips":["Always create swizzles from the curated atom constants, never hand-written triples","Centralize swizzle selection in one config helper shared by layout construction and descriptor building"],"tags":["cuda","cutlass","cute-dsl","umma","smem-descriptor","swizzle","blackwell"],"backgroundTag":"invalid-swizzle-mode","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}