{"record":{"id":"f6607b660959cbbc","repo":"jax-ml/jax","slug":"unknown-core-type-core-type","errorCode":null,"errorMessage":"Unknown core type: {core_type}","messagePattern":"Unknown core type: (.+?)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/tpu_custom_call.py","lineNumber":546,"sourceCode":"\n  def assign_device_type_based_on_core_type(op: ir.Operation) -> ir.WalkResult:\n    nonlocal sparsecore_func_found\n    nonlocal tensorcore_func_found\n    if op.name == \"func.func\":\n      if \"tpu.core_type\" in op.attributes:\n        core_type = op.attributes[\"tpu.core_type\"]\n        if str(core_type) in [\n            f\"#tpu.core_type<{c}>\"\n            for c in [\"sc_scalar_subcore\", \"sc_vector_subcore\"]\n        ]:\n          sparsecore_func_found = True\n          if tensorcore_func_found:\n            return ir.WalkResult.INTERRUPT\n          return ir.WalkResult.SKIP\n        if str(core_type) == \"#tpu.core_type<tc>\":\n          tensorcore_func_found = True\n          return ir.WalkResult.SKIP\n        raise ValueError(f\"Unknown core type: {core_type}\")\n    return ir.WalkResult.ADVANCE\n\n  module.operation.walk(\n      assign_device_type_based_on_core_type, walk_order=ir.WalkOrder.PRE_ORDER\n  )\n  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(","sourceCodeStart":528,"sourceCodeEnd":564,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/tpu_custom_call.py#L528-L564","documentation":"While scanning the lowered Mosaic module, an unknown #tpu.core_type attribute was found — only <tc> (tensor core) and the sparsecore variant are recognized.","triggerScenarios":"A lowered Pallas module containing a core_type attribute other than the supported tc/sparsecore forms, often from a version mismatch between Mosaic/Pallas IR and this JAX.","commonSituations":"Mixing JAX version with a newer/older Pallas or Mosaic that emits new core types (e.g. new chip types).","solutions":["Align JAX, Pallas, and Mosaic/jaxlib versions","Re-lower the kernel with the toolchain matching your JAX version"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"import jax, jaxlib\nassert version_parse(jax.__version__) >= version_parse(min_required)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Pin jax/jaxlib/pallas versions together","Never reuse lowered Mosaic asm across JAX upgrades"],"tags":["tpu","pallas","core-type","version-mismatch","mlir"],"backgroundTag":"unsupported-enum-value","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}