{"record":{"id":"0a88d1432462e8be","repo":"jax-ml/jax","slug":"bkv-must-be-a-multiple-of-num-lanes","errorCode":null,"errorMessage":"{bkv=} must be a multiple of {NUM_LANES}.","messagePattern":"(.+?) must be a multiple of (.+?)\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/pallas/ops/tpu/splash_attention/splash_attention_kernel.py","lineNumber":1465,"sourceCode":"        f\"{bq=} should not be greater than {q_seq_len=}\")\n  if bkv > kv_seq_len:\n    raise ValueError(\n        f\"{bkv=} should not be greater than {kv_seq_len=}\")\n\n  if not is_mqa and num_q_heads % num_kv_heads != 0:\n    raise ValueError(\n        f\"In MHA, expected number of 'key' heads ({num_kv_heads}) to be a\"\n        f\" multiple of the number of 'query' heads ({num_q_heads})\"\n    )\n\n  if k.shape[:-1] != v.shape[:-1]:\n    raise ValueError(\n        f\"Expected 'key' {k.shape} and 'value' {v.shape} to have the same \"\n        \"leading dimensions.\"\n    )\n\n  if bkv % NUM_LANES:\n    raise ValueError(f\"{bkv=} must be a multiple of {NUM_LANES}.\")\n\n  # TODO(amagni/sharadmv): when adding block_compute, make sure that is a\n  # multiple of NUM_LANES.\n\n  q_heads_per_kv_head = num_q_heads // num_kv_heads\n\n  if mask_info.data_next is not None:\n    grid_width = mask_info.data_next.shape[-1]\n  else:\n    grid_width = kv_seq_len // bkv\n\n  grid = (num_q_heads, q_seq_len // bq, grid_width)\n\n  def o_index_map(h, i, *_):\n    return h, i, 0\n\n  o_spec = pl.BlockSpec((None, bq, head_dim_v), o_index_map)\n","sourceCodeStart":1447,"sourceCodeEnd":1483,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/pallas/ops/tpu/splash_attention/splash_attention_kernel.py#L1447-L1483","documentation":"In the backward dq kernel, block_kv must be divisible by NUM_LANES=8 to keep lane-aligned softmax reductions correct during the gradient computation. The source notes block_kv_compute alignment is a TODO, but bkv itself is enforced.","triggerScenarios":"jax.grad through splash attention with block_sizes.block_kv=100 or any value not divisible by 8 (e.g. tuned for HBM size 7*16=112).","commonSituations":"Hand-tuned block sizes that satisfy the forward constraints but forgot the backward's NUM_LANES constraint on block_kv.","solutions":["Set block_kv to a multiple of 8 (64, 128, 256)","Use default block sizes (block_sizes=None) which are lane-aligned","Validate all block sizes against NUM_LANES=8 in your config loader"],"exampleFix":"// before\nBlockSizes(block_kv=112)\n// after\nBlockSizes(block_kv=128)","handlingStrategy":"validation","validationCode":"assert block_sizes.block_kv % 8 == 0","typeGuard":"def lane_aligned(x: int) -> bool: return x % 8 == 0","tryCatchPattern":null,"preventionTips":["Validate backward-path block sizes too, not just forward","Keep block sizes powers of two"],"tags":["jax","pallas","tpu","splash-attention","block-size","backward"],"backgroundTag":"block-size-constraint-violation","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}