{"record":{"id":"1c7978969953254c","repo":"jax-ml/jax","slug":"the-current-tpu-does-not-have-sparsecores","errorCode":null,"errorMessage":"The current TPU does not have SparseCores","messagePattern":"The current TPU does not have SparseCores","errorType":"exception","errorClass":"RuntimeError","httpStatus":null,"severity":"critical","filePath":"jax/_src/pallas/mosaic/sc_core.py","lineNumber":41,"sourceCode":"\nimport jax\nfrom jax._src import core as jax_core\nfrom jax._src import tree_util\nfrom jax._src.pallas import core as pallas_core\nfrom jax._src.pallas.mosaic import core as tpu_core\nfrom jax._src.pallas.mosaic import tpu_info\nimport jax.numpy as jnp\n\n\ndef get_sparse_core_info() -> tpu_info.SparseCoreInfo:\n  \"\"\"Returns the SparseCore information for the current device.\n\n  Raises:\n    RuntimeError: If the current TPU does not have SparseCores.\n  \"\"\"\n  sc_info = tpu_info.get_tpu_info().sparse_core\n  if sc_info is None:\n    raise RuntimeError(\"The current TPU does not have SparseCores\")\n  return sc_info\n\n\n@dataclasses.dataclass(frozen=True, kw_only=True)\nclass ScalarSubcoreMesh(pallas_core.Mesh):\n  axis_name: str\n  num_cores: int = dataclasses.field(\n      default_factory=lambda: get_sparse_core_info().num_cores\n  )\n\n  def __post_init__(self):\n    sc_info = get_sparse_core_info()\n    if self.num_cores > sc_info.num_cores:\n      raise ValueError(\n          f\"Mesh has {self.num_cores} cores, but the current TPU chip has only\"\n          f\" {sc_info.num_cores} SparseCores\"\n      )\n","sourceCodeStart":23,"sourceCodeEnd":59,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic/sc_core.py#L23-L59","documentation":"get_sparse_core_info queries the current TPU's SparseCore metadata; on chips without SparseCores (e.g. TPU v2/v3 or CPU/localhost) sparse_core is None and a RuntimeError is raised. Most SC-core primitives and mesh classes call this at construction time.","triggerScenarios":"Constructing ScalarSubcoreMesh/VectorSubcoreMesh, calling supported_shapes, or using SC-core lowering rules on hardware without SparseCores, or calling tpu_info on a non-TPU host.","commonSituations":"Running a SparseCore Pallas kernel locally or in CI on CPU; targeting TPU v4+ only code on older TPU generations; using the wrong backend (gpu/cpu instead of tpu).","solutions":["Run on TPU hardware that has SparseCores (v4/v5p/v6e etc.) by selecting jax.devices('tpu')","Guard the mesh construction with get_sparse_core_info() in a try/except and skip SC paths when unavailable","In tests/CI, gate SparseCore tests on device availability"],"exampleFix":"# before\nmesh = VectorSubcoreMesh(...)\n# after\ntry:\n  info = get_sparse_core_info()\nexcept RuntimeError:\n  info = None\nif info is not None:\n  mesh = VectorSubcoreMesh(...)","handlingStrategy":"try-catch","validationCode":"import jax\ndevs = jax.devices()\nhas_sc = any(d.platform == 'tpu' for d in devs)\n# and expect tpu_info.sparse_core to be non-None on SC-capable chips","typeGuard":"def sparse_cores_available() -> bool:\n    try:\n        get_sparse_core_info()\n        return True\n    except RuntimeError:\n        return False","tryCatchPattern":"try:\n    sc_info = get_sparse_core_info()\nexcept RuntimeError as e:\n    if 'SparseCores' in str(e):\n        raise SystemExit('This kernel requires a SparseCore-capable TPU') from e\n    raise","preventionTips":["Gate SparseCore code on device/platform checks","Run SC tests only on SC-capable TPU backends"],"tags":["jax","pallas","tpu","sparsecore","hardware","environment"],"backgroundTag":"unsupported-hardware","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}