{"record":{"id":"222f1888a7b7b3ee","repo":"jax-ml/jax","slug":"only-single-1-and-double-2-buffering-are-suppo","errorCode":null,"errorMessage":"Only single (1) and double (2) buffering are supported. Got {buffer_count}.","messagePattern":"Only single \\(1\\) and double \\(2\\) buffering are supported\\. Got (.+?)\\.","errorType":"validation","errorClass":"LoweringException","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic/lowering.py","lineNumber":1276,"sourceCode":"            bd.padding if isinstance(bd, pallas_core.Element) else (0, 0)\n            for bd in bm.block_shape\n        ]\n        pad_low, pad_high = map(list, zip(*padding))\n        block_params[\"window_kind\"] = ir.Attribute.parse(\n            f\"#tpu.element_window<{pad_low},{pad_high}>\"\n        )\n      if pipeline_mode is not None:\n        if not isinstance(pipeline_mode, pallas_core.Buffered):\n          raise LoweringException(\n              f\"Unsupported pipeline mode: {pipeline_mode}.\"\n          )\n        if pipeline_mode.use_lookahead:\n          raise NotImplementedError(\n              \"Lookahead is not supported for XLA pipeline emitter lowering.\"\n          )\n        buffer_count = pipeline_mode.buffer_count\n        if buffer_count < 1 or buffer_count > 2:\n          raise LoweringException(\n              \"Only single (1) and double (2) buffering are supported. Got\"\n              f\" {buffer_count}.\"\n          )\n        pipeline_mode_str = \"synchronous\" if buffer_count == 1 else \"double_buffered\"\n        block_params[\"pipeline_mode\"] = ir.Attribute.parse(\n            f\"#tpu.pipeline_mode<{pipeline_mode_str}>\"\n        )\n        if pipeline_mode.revisit is not None:\n          if (\n              pipeline_mode.revisit == pallas_core.RevisitMode.ANY\n              and buffer_count > 1\n          ):\n            raise LoweringException(\n                \"RevisitMode.ANY is not supported with double buffering.\"\n            )\n          block_params[\"revisit_mode\"] = ir.Attribute.parse(\n              f\"#tpu.revisit_mode<{pipeline_mode.revisit.value}>\"\n          )","sourceCodeStart":1258,"sourceCodeEnd":1294,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic/lowering.py#L1258-L1294","documentation":"The TPU lowering supports only single buffering (buffer_count=1, synchronous) and double buffering (buffer_count=2) for pipelined DMA. Any other buffer_count (0, 3, ...) raises LoweringException, because the emitted #tpu.pipeline_mode attribute can only be 'synchronous' or 'double_buffered'.","triggerScenarios":"Passing pipeline_mode=pallas_core.Buffered(buffer_count=3) (or 0) to the TPU lowering; requesting triple buffering for deeper DMA overlap.","commonSituations":"Tuning kernels by increasing buffer counts expecting more prefetch depth; porting GPU-style multi-stage pipelining to Pallas TPU; scripts parameterizing buffer_count without clamping.","solutions":["Use buffer_count=2 (double buffering), the maximum supported","Use buffer_count=1 for synchronous, memory-frugal pipelines","Restructure the kernel (e.g. smaller tiles) instead of adding more buffers to hide latency"],"exampleFix":"# before\npipeline_mode=pallas_core.Buffered(buffer_count=3)\n\n# after\npipeline_mode=pallas_core.Buffered(buffer_count=2)","handlingStrategy":"validation","validationCode":"from jax._src.pallas import pallas_core\ndef check_buffer_count(pm):\n    if isinstance(pm, pallas_core.Buffered) and not (1 <= pm.buffer_count <= 2):\n        raise ValueError('buffer_count must be 1 or 2 on TPU')","typeGuard":"def has_supported_buffer_count(pm) -> bool:\n    from jax._src.pallas import pallas_core\n    return not isinstance(pm, pallas_core.Buffered) or pm.buffer_count in (1, 2)","tryCatchPattern":"try:\n    lower_jaxpr_into_pipelined_module(..., pipeline_mode=pm)\nexcept LoweringException as e:\n    if 'buffering' in str(e):\n        pm = pallas_core.Buffered(buffer_count=min(pm.buffer_count, 2)); retry","preventionTips":["Clamp buffer_count to 1 or 2 before passing pipeline_mode","Use smaller tiles rather than more buffers to tune performance"],"tags":["jax","pallas","tpu","pipelining","buffer-count"],"backgroundTag":"pallas-unsupported-buffer-count","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}