{"record":{"id":"50ff1ddd6ef63287","repo":"xai-org/x-algorithm","slug":"m-must-be-64-128-or-256","errorCode":null,"errorMessage":"M must be 64, 128 or 256","messagePattern":"M must be 64, 128 or 256","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":164,"sourceCode":"    a_type,\n    b_type,\n    c_type,\n    M: int,\n    N: int,\n    a_major: Major,\n    b_major: Major,\n    a_neg: ScaleIn = ScaleIn.One,\n    b_neg: ScaleIn = ScaleIn.One,\n    c_sat: Saturate = Saturate.False_,\n    is_sparse: bool = False,\n    max_shift: MaxShift = MaxShift.NoShift,\n) -> int:\n    a_fmt = int(to_UMMA_format(a_type))\n    b_fmt = int(to_UMMA_format(b_type))\n    c_fmt = int(to_C_format(c_type))\n\n    if M not in (64, 128, 256):\n        raise ValueError(\"M must be 64, 128 or 256\")\n    if N < 8 or N > 256 or (N & 7):\n        raise ValueError(\"N must be a multiple of 8 in the range 8…256\")\n\n    m_dim = M >> 4\n    n_dim = N >> 3\n\n    desc = 0\n    desc |= (0 & 0x3) << 0\n    desc |= (int(is_sparse) & 0x1) << 2\n    desc |= (int(c_sat) & 0x1) << 3\n    desc |= (c_fmt & 0x3) << 4\n    desc |= (a_fmt & 0x7) << 7\n    desc |= (b_fmt & 0x7) << 10\n    desc |= (int(a_neg) & 0x1) << 13\n    desc |= (int(b_neg) & 0x1) << 14\n    desc |= (int(a_major) & 0x1) << 15\n    desc |= (int(b_major) & 0x1) << 16\n    desc |= (n_dim & 0x3F) << 17","sourceCodeStart":146,"sourceCodeEnd":182,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L146-L182","documentation":"make_instr_desc encodes the UMMA instruction M dimension into the instruction descriptor and hardware only supports M in {64, 128, 256}. Passing any other M (e.g. 32, 96, 512) raises ValueError before the descriptor is built.","triggerScenarios":"Constructing an MMA op / calling mma_op_to_idesc where the tile M is derived from a head count or block size like 32, 96, or 512, which is not in (64, 128, 256).","commonSituations":"Tuning attention tile shapes (e.g. 96-row Q tiles for small batch) or porting SM90 tile configs to SM100 where the legal M set differs.","solutions":["Round the MMA M dimension to 64, 128, or 256 (e.g. pad a 96-row tile to 128)","Update the tile-shape config so the atom's M is legal while padding handles the remainder at the layout level","If a smaller M is required, use masking over a 64-row atom rather than an unsupported atom size"],"exampleFix":"# before\ndesc = make_instr_desc(a, b, c, M=96, N=128)\n# after\ndesc = make_instr_desc(a, b, c, M=128, N=128)  # mask remainder rows","handlingStrategy":"validation","validationCode":"if M not in (64, 128, 256):\n    raise ValueError(f'M={M} illegal for SM100 UMMA; use 64/128/256')","typeGuard":"def is_legal_umma_m(m: int) -> bool:\n    return m in (64, 128, 256)","tryCatchPattern":null,"preventionTips":["Validate tile shapes once at autotune/setup time, not per launch","Pad problem sizes to legal atoms and mask, rather than inventing illegal atom sizes"],"tags":["cutlass","tcgen05","mma","tile-shape","validation"],"backgroundTag":"invalid-argument-value","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}