{"record":{"id":"0d7f09fc808f6765","repo":"xai-org/x-algorithm","slug":"swizzle-128b-base32b-is-invalid-for-major-k","errorCode":null,"errorMessage":"SWIZZLE_128B_BASE32B is invalid for Major-K","messagePattern":"SWIZZLE_128B_BASE32B is invalid for Major-K","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":262,"sourceCode":"    if major is Major.MN:\n        swizzle_atom_k_size = 4 if layout_type is LayoutType.SWIZZLE_128B_BASE32B else 8\n        canonical_layout = cute.logical_divide(layout, (swizzle_atom_mn_size, swizzle_atom_k_size))\n        if not cute.is_congruent(canonical_layout, ((1, 1), (1, 1))):\n            raise ValueError(\"Not a canonical UMMA_MN Layout: Expected profile failure.\")\n        stride_00 = canonical_layout.stride[0][0]\n        if layout_type is not LayoutType.SWIZZLE_NONE and stride_00 != 1:\n            raise ValueError(\"Not a canonical UMMA_MN Layout: Expected stride failure.\")\n        stride_10 = canonical_layout.stride[1][0]\n        if stride_10 != swizzle_atom_mn_size:\n            raise ValueError(\"Not a canonical UMMA_MN Layout: Expected stride failure.\")\n        stride_01, stride_11 = canonical_layout.stride[0][1], canonical_layout.stride[1][1]\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(\"Not a canonical UMMA_K Layout: Expected MN-size multiple of 8.\")\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    desc = 0\n    desc |= (leading_byte_offset & 0x3FFF) << 16\n    desc |= (stride_byte_offset & 0x3FFF) << 32\n    desc |= (VERSION & 0x3) << 46","sourceCodeStart":244,"sourceCodeEnd":280,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L244-L280","documentation":"The 128B_BASE32B swizzle mode (Swizzle<2,5,2>, used for 32-bit-base fp8/mxfp8 tcgen05 operands on Blackwell) has no descriptor encoding for K-major operands in this builder — the PTX smem descriptor format only defines it for MN-major layouts. make_smem_desc_base rejects the combination up front.","triggerScenarios":"Calling make_smem_desc_base (directly or via gemm_ptx*/smem_desc_base_from_tensor) with swizzle Swizzle<2,5,2> and major=Major.K, e.g. an fp8 B-operand stored K-major in shared memory.","commonSituations":"Using mxfp8/fp8 tensors where the B operand arrives K-major from a transpose-free GEMM epilogue; porting a CUTLASS CollectiveBuilder config with K-major B and reusing the 128B_BASE32B smem atom.","solutions":["Store the operand MN-major (swap to an MN-major smem layout) so Swizzle<2,5,2> is legal","If K-major is required, switch to a standard swizzle (Swizzle<3,4,3> 128B or <2,4,3> 64B) instead of 128B_BASE32B","Set the MMA operand major mode to MN in the MmaOp so mma_op_to_idesc and the smem descriptor agree"],"exampleFix":"# before\ndesc = make_smem_desc_base(kmajor_layout, cute.Swizzle(2,5,2), Major.K)\n# after\ndesc = make_smem_desc_base(mnmajor_layout, cute.Swizzle(2,5,2), Major.MN)","handlingStrategy":"type-guard","validationCode":"if sw.num_base == 5:  # Swizzle<2,5,2> = 128B_BASE32B\n    assert major is Major.MN, '128B_BASE32B requires MN-major operand'","typeGuard":"def base32b_ok(sw: cute.Swizzle, major: Major) -> bool:\n    return not (sw.num_base == 5 and major is Major.K)","tryCatchPattern":null,"preventionTips":["For fp8/mxfp8 operands plan the B matrix as MN-major from the start","Keep a table of swizzle x allowed-major modes near kernel config"],"tags":["cuda","cutlass","cute-dsl","umma","fp8","blackwell","smem-descriptor"],"backgroundTag":"invalid-major-mode","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}