{"record":{"id":"3c1352a3f4217e6d","repo":"jax-ml/jax","slug":"the-iteration-bounds-and-dimension-semantics-attri","errorCode":null,"errorMessage":"The iteration bounds and dimension semantics attributes must have the same number of elements.","messagePattern":"The iteration bounds and dimension semantics attributes must have the same number of elements\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/tpu_custom_call.py","lineNumber":570,"sourceCode":"  if tensorcore_func_found and sparsecore_func_found:\n    raise ValueError(\n        \"A single Mosaic kernel cannot contain both TensorCore and SparseCore\"\n        \" functions.\"\n    )\n  if sparsecore_func_found:\n    return \"sparsecore\"\n  return None\n\n\ndef _get_active_core_count(module: ir.Module) -> int | None:\n\n  def get_core_parallel_dim_size(\n      dim_semantics: ir.ArrayAttr,\n      iter_bounds: ir.DenseI64ArrayAttr,\n      other_subkernel_core_dim_size: int | None = None) -> int | None:\n\n    if len(iter_bounds) != len(dim_semantics):\n      raise ValueError(\n          \"The iteration bounds and dimension semantics attributes must have\"\n          \" the same number of elements.\"\n      )\n\n    subkernel_core_dim_size = None\n\n    for dim_idx, (dim_size, dim_sem) in enumerate(\n        zip(iter_bounds, dim_semantics)\n    ):\n      if str(dim_sem) != \"#tpu.dimension_semantics<core_parallel>\":\n        continue\n\n      if ir.ShapedType.is_dynamic_size(dim_size):\n        raise ValueError(\n            \"The iteration bound corresponding to the core-parallel dimension \"\n            f\"{dim_idx} must be statically known.\"\n        )\n      if subkernel_core_dim_size is not None:","sourceCodeStart":552,"sourceCodeEnd":588,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/tpu_custom_call.py#L552-L588","documentation":"A Mosaic subkernel's iteration bounds array and dimension semantics array have different lengths; they must be parallel 1:1.","triggerScenarios":"A lowered tpu.subkernel op whose in_bounds length differs from dim_semantics length — typically from malformed/incompatible Mosaic IR.","commonSituations":"Version mismatch between Mosaic emitter and JAX lowering; manually constructed IR.","solutions":["Update JAX/jaxlib to match the Mosaic version that produced the module","Regenerate the kernel rather than reusing stale lowered IR"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"assert len(in_bounds) == len(dim_semantics) for each subkernel before lowering","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Regenerate lowered IR with the matching toolchain","Don't hand-edit subkernel attributes"],"tags":["tpu","pallas","mlir","subkernel","validation"],"backgroundTag":"schema-validation-failed","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}