{"record":{"id":"cfa9230266e8e63d","repo":"sgl-project/sglang","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":"python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py","lineNumber":221,"sourceCode":"\n    if M == 4:  # Swizzle<*,4,3>\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        }[\n            B\n        ]  # KeyError ⇒ invalid B→ raise\n    if M == 5:  # Swizzle<2,5,2> (the only legal triple for 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    # Any other (M,B,S) triple is not a UMMA-legal shared-memory layout\n    raise ValueError(\"Unsupported swizzle triple for UMMA smem descriptor\")\n\n\ndef make_smem_desc_base(\n    layout: cute.Layout, swizzle: cute.Swizzle, major: Major\n) -> int:\n    \"\"\"\n    Convert a 2-D *shared-memory* Cute layout into the Blackwell 64-bit\n    smem-descriptor, without the smem start address.\n    layout must correspond to layout of an uint128 tensor.\n    \"\"\"\n    # ------------------------------------------------------------------ meta\n    layout_type = _layout_type(swizzle)  # resolve SWIZZLE_* family\n\n    VERSION = 1  # bits 46–47\n    LBO_MODE = 0  # bit  52\n    BASE_OFFSET = 0  # bits 49–51   (CUTLASS always 0)\n\n    # ---------------------------------------------------------- strides  (units: uint128_t = 16 B)","sourceCodeStart":203,"sourceCodeEnd":239,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py#L203-L239","documentation":"_layout_type raises when the swizzle triple (num_base, num_bits, num_shift) matches neither the M==4 (Swizzle<*,4,3>) nor M==5 (Swizzle<2,5,2>) families — i.e. it is not a UMMA-legal shared-memory layout at all.","triggerScenarios":"Calling make_smem_desc_base with arbitrary swizzles such as Swizzle<1,3,0>, Swizzle<3,4,3>, or a non-swizzled layout's default swizzle object with unusual fields.","commonSituations":"Deriving swizzle parameters programmatically from tensor shapes/strides and landing on an unsupported combination; using SM90 wgmma-era swizzles with the SM100 UMMA descriptor builder.","solutions":["Use LayoutType.SWIZZLE_NONE-compatible layouts (no swizzle) or one of the legal triples: Swizzle<0..3,4,3> or Swizzle<2,5,2>.","Copy the swizzle from a working kernel's atom in this same directory rather than computing one.","Consult the UMMA smem descriptor table (32B/64B/128B swizzle atoms) when choosing layouts."],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":null,"typeGuard":"def is_umma_legal_swizzle(sw):\n    B, M, S = sw.num_bits, sw.num_base, sw.num_shift\n    return (M == 4 and S == 3 and 0 <= B <= 3) or (M == 5 and (B, S) == (2, 2))","tryCatchPattern":null,"preventionTips":["Restrict swizzles to the two legal families (Swizzle<0..3,4,3>, Swizzle<2,5,2>).","Don't port SM90 wgmma swizzles into SM100 UMMA descriptor code."],"tags":["cutlass","sm100","swizzle","shared-memory"],"backgroundTag":"invalid-swizzle-layout","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}