{"record":{"id":"ff80cef2ce28da1b","repo":"jax-ml/jax","slug":"m-must-be-a-multiple-of-16-and-256-got-m","errorCode":null,"errorMessage":"M must be a multiple of 16 and <= 256, got: {m}","messagePattern":"M must be a multiple of 16 and <= 256, got: (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":108,"sourceCode":"      return 1\n    elif ty == ir.IntegerType.get_signless(8):  # Only s8 for now.\n      assert acc_dtype == i32\n      return 1\n    else:\n      raise NotImplementedError(f\"Unsupported input dtype: {ty}\")\n  a_type_val = get_input_encoding(a_dtype)\n  b_type_val = get_input_encoding(b_dtype)\n  desc |= (a_type_val << 7)   # A dtype, bits 7-9\n  desc |= (b_type_val << 10)  # B dtype, bits 10-12\n  # We ignore negate bits 13-14\n  desc |= transpose_a << 15  # Transpose A\n  desc |= transpose_b << 16  # Transpose B\n  if n % 8 or n > 256:\n    raise ValueError(f\"N must be a multiple of 8 and <= 256, got: {n}\")\n  desc |= (n >> 3) << 17  # N, bits 17-22\n  # Bit 23 is reserved\n  if m % 16 or m > 256:\n    raise ValueError(f\"M must be a multiple of 16 and <= 256, got: {m}\")\n  desc |= (m >> 4) << 24  # M >> 4, bits 24-28\n  # Bit 29 is reserved\n  # We ignore max shift under .ws, bits 30-31\n  return arith.constant(ir.IntegerType.get_signless(32), desc)\n\n\ndef _create_scaled_instr_descriptor(\n    get_input_encoding: Callable[[ir.Type], int],\n    m: int,\n    n: int,\n    a_type: ir.Type,\n    b_type: ir.Type,\n    a_scale_idx: int,\n    b_scale_idx: int,\n    transpose_a: bool,\n    transpose_b: bool,\n    scale_type: ir.Type,\n    sparse: bool = False,","sourceCodeStart":90,"sourceCodeEnd":126,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L90-L126","documentation":"M is packed into descriptor bits 24-28 as m>>4, so M must be a multiple of 16 and no larger than 256. create_instr_descriptor rejects other M values with ValueError.","triggerScenarios":"Calling tcgen05 mma with M=100 (not multiple of 16) or M=512 (> 256).","commonSituations":"Tile-size tuning sweeps that step M by 8; assuming M can be as large as N.","solutions":["Use M in {16, 32, ..., 256}","Split larger M across multiple MMA instructions/tiles"],"exampleFix":"# before\nmma(acc, a, b, m=264, n=128)\n\n# after\n# split into two m=132? no — use m=128 twice\nmma(acc1, a1, b, m=128, n=128)\nmma(acc2, a2, b, m=128, n=128)","handlingStrategy":"validation","validationCode":"assert m % 16 == 0 and m <= 256, f'M={m} must be multiple of 16 and <= 256'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Use M tiles from {16..256 step 16}","Split M > 256 across multiple mma calls"],"tags":["tcgen05","matmul","mosaic","gpu","shape-validation"],"backgroundTag":"matrix-dimension-constraint","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}