{"record":{"id":"ed53ba29e6e6957b","repo":"jax-ml/jax","slug":"tpu-version-must-be-4-or-higher","errorCode":null,"errorMessage":"TPU version must be 4 or higher.","messagePattern":"TPU version must be 4 or higher\\.","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/experimental/pallas/ops/tpu/ragged_paged_attention/tuned_block_sizes.py","lineNumber":1458,"sourceCode":"  if num_devices is not None:\n    name += f'-{num_devices}'\n  return name\n\n\ndef get_tuned_block_sizes(\n    q_dtype,\n    kv_dtype,\n    num_q_heads_per_blk,\n    num_kv_heads_per_blk,\n    head_dim,\n    page_size,\n    max_num_batched_tokens,\n    pages_per_seq,\n) -> tuple[int, int]:\n  \"\"\"Look up for the best (num_kv_pages_per_blk, num_queries_per_blk) from auto-tuned table.\"\"\"\n  tpu_version = get_tpu_version()\n  if tpu_version < 4:\n    raise NotImplementedError('TPU version must be 4 or higher.')\n  key = (\n      q_dtype,\n      kv_dtype,\n      num_q_heads_per_blk,\n      num_kv_heads_per_blk,\n      head_dim,\n      page_size,\n      max_num_batched_tokens,\n      pages_per_seq,\n  )\n  key = simplify_key(key)\n  device_name = get_device_name()\n\n  # Default block sizes.\n  bkv, bq = (128, 32)\n  if tpu_version == 4:\n    # This default block size is not tuned, only make sure there's no\n    # OOM in vmem","sourceCodeStart":1440,"sourceCodeEnd":1476,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/pallas/ops/tpu/ragged_paged_attention/tuned_block_sizes.py#L1440-L1476","documentation":"get_tuned_block_sizes looks up auto-tuned (num_kv_pages_per_blk, num_queries_per_blk) values for the TPU ragged paged attention kernel; the tuning tables only cover TPU v4 and later. On older TPUs (v2/v3) or when the TPU version is detected as < 4 it raises NotImplementedError.","triggerScenarios":"Running ragged_paged_attention with automatic block-size tuning on TPU v2/v3 hardware, or on a runtime/SDK where get_tpu_version() misdetects or returns a low version.","commonSituations":"Deploying JAX code written for v4/v5p/v6e onto older Cloud TPU slices or an emulator; running on CPU/GPU while the pallas TPU path is still imported and invoked.","solutions":["Run on TPU v4 or newer (v4, v5e, v5p, v6e)","Explicitly pass num_kv_pages_per_block and num_queries_per_block if the API path allows bypassing the tuned table","Check jax.devices() to confirm you are actually attached to a TPU and its version","Update libtpu/JAX so TPU version detection is correct"],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":"import jax\ndevs = jax.devices()\nassert all(d.platform == 'tpu' for d in devs), 'requires TPU'","typeGuard":"def is_tpu_v4_plus() -> bool:\n    try:\n        from jax.experimental.pallas.ops.tpu.ragged_paged_attention.tuned_block_sizes import get_tpu_version\n        return get_tpu_version() >= 4\n    except Exception:\n        return False","tryCatchPattern":"try:\n    sizes = get_tuned_block_sizes(...)\nexcept NotImplementedError:\n    sizes = None  # fall back to default block sizes","preventionTips":["Gate TPU-kernel usage on device platform and version at startup","Pass explicit block sizes when tuning tables are unavailable"],"tags":["jax","pallas","tpu","hardware-version","not-implemented"],"backgroundTag":"unsupported-tpu-version","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}