{"record":{"id":"59b3f80b416686ce","repo":"jax-ml/jax","slug":"n-must-be-a-multiple-of-8-and-256-got-n","errorCode":null,"errorMessage":"N must be a multiple of 8 and <= 256, got: {n}","messagePattern":"N must be a multiple of 8 and <= 256, got: (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":104,"sourceCode":"      assert acc_dtype in {f16, f32}\n      return 0\n    elif ty == ir.Float8E5M2Type.get():\n      assert acc_dtype in {f16, f32}\n      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,","sourceCodeStart":86,"sourceCodeEnd":122,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L86-L122","documentation":"The tcgen05 MMA instruction descriptor packs N into bits 17-22 as n>>3, requiring N to be a multiple of 8 and at most 256. Violating either makes the descriptor unrepresentable, so create_instr_descriptor raises ValueError.","triggerScenarios":"Calling the tcgen05 mma path with N like 100 (not multiple of 8) or 264 (> 256).","commonSituations":"Using arbitrary matrix shapes from a model config; padding shapes to 128/256 everywhere except N.","solutions":["Pad N up to the next multiple of 8 (typically 64/128/256)","Choose tile sizes from supported set: 8, 16, ..., 256","Validate shapes before building the descriptor"],"exampleFix":"# before\nmma(acc, a, b, m=128, n=100)\n\n# after\nmma(acc, a, b, m=128, n=104)  # pad to multiple of 8","handlingStrategy":"validation","validationCode":"assert n % 8 == 0 and n <= 256, f'N={n} must be multiple of 8 and <= 256'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Restrict tile sizes to multiples of 8 (prefer 64/128/256)","Validate N before building the instruction descriptor"],"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"}