{"record":{"id":"bfd4a1c8b6a16bf6","repo":"jax-ml/jax","slug":"a-single-mosaic-subkernel-cannot-contain-multiple","errorCode":null,"errorMessage":"A single Mosaic subkernel cannot contain multiple core sharding dimensions.","messagePattern":"A single Mosaic subkernel cannot contain multiple core sharding dimensions\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/tpu_custom_call.py","lineNumber":589,"sourceCode":"          \"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:\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:","sourceCodeStart":571,"sourceCodeEnd":607,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/tpu_custom_call.py#L571-L607","documentation":"A subkernel declares more than one core_parallel dimension; only a single core-sharding dimension per subkernel is allowed.","triggerScenarios":"Annotating two dimensions as #tpu.dimension_semantics<core_parallel> in the same subkernel.","commonSituations":"Hand-written or migrated Mosaic kernels that shard across two core axes.","solutions":["Keep only one core_parallel dimension; move other dims to arith/layout semantics"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"assert sum(str(d) == '#tpu.dimension_semantics<core_parallel>' for d in dim_semantics) <= 1","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Annotate exactly one core-parallel dim per subkernel"],"tags":["tpu","pallas","core-parallel","subkernel","validation"],"backgroundTag":"unsupported-configuration","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}