{"record":{"id":"b26951bf0bdab704","repo":"jax-ml/jax","slug":"the-iteration-bound-corresponding-to-the-core-para-b26951","errorCode":null,"errorMessage":"The iteration bound corresponding to the core-parallel dimension be the same across all subkernels.","messagePattern":"The iteration bound corresponding to the core-parallel dimension be the same across all subkernels\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/tpu_custom_call.py","lineNumber":597,"sourceCode":"    ):\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:\n        raise ValueError(\n            \"A single Mosaic subkernel cannot contain multiple core sharding \"\n            \"dimensions.\"\n        )\n      if (\n          other_subkernel_core_dim_size is not None\n          and other_subkernel_core_dim_size != dim_size\n      ):\n        raise ValueError(\n            \"The iteration bound corresponding to the core-parallel dimension \"\n            \"be the same across all subkernels.\"\n        )\n      subkernel_core_dim_size = dim_size\n\n    return subkernel_core_dim_size\n\n  core_parallel_dim_size = None\n\n  for op in module.body.operations:\n    if op.operation.name != \"func.func\":\n      continue\n\n    if (\n        \"iteration_bounds\" not in op.attributes\n        or \"dimension_semantics\" not in op.attributes\n    ):\n      continue","sourceCodeStart":579,"sourceCodeEnd":615,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/tpu_custom_call.py#L579-L615","documentation":"Different subkernels in the same Mosaic kernel declare core_parallel iteration bounds with different static sizes; all subkernels must agree.","triggerScenarios":"Composite kernels where one subkernel uses core dim size N and another uses M != N.","commonSituations":"Fusing subkernels written for different core counts.","solutions":["Make the core-parallel bound identical across all subkernels","Recompile subkernels for the same core count"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"sizes = {get_core_parallel_dim_size(sk) for sk in subkernels}\nassert len(sizes) <= 1","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Parameterize core count once and reuse across subkernels"],"tags":["tpu","pallas","core-parallel","subkernel","consistency"],"backgroundTag":"inconsistent-configuration","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}