{"record":{"id":"076e9d3faed872c9","repo":"jax-ml/jax","slug":"only-positional-arguments-are-supported-by-debug-p-076e9d","errorCode":null,"errorMessage":"Only positional arguments are supported by debug_print on Pallas.","messagePattern":"Only positional arguments are supported by debug_print on Pallas\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/triton/lowering.py","lineNumber":1332,"sourceCode":"    fmt: str,\n    ordered,\n    partitioned,\n    in_tree,\n    static_args,\n    np_printoptions,\n    has_placeholders,\n    logging_record,\n):\n  del partitioned, np_printoptions\n  if ordered:\n    raise NotImplementedError(\"Ordered debug_print is not supported on Pallas.\")\n  if has_placeholders:\n    raise ValueError(\n        \"pl.debug_print() does not support placeholders when lowering to Triton\"\n    )\n  args, kwargs = debugging.merge_callback_args(in_tree, args, static_args)\n  if kwargs:\n    raise ValueError(\n        \"Only positional arguments are supported by debug_print on Pallas.\"\n    )\n\n  tt_dialect.print_(\n      f\" {fmt} \",\n      hex=False,\n      args=args,\n      is_signed=ir.DenseI32ArrayAttr.get([\n          jnp.issubdtype(aval.dtype, jnp.signedinteger) for aval in ctx.avals_in\n      ]),\n  )\n  return ()\n\n\ndef _set_attr(v: ir.Value, name: str, attr: ir.Attribute) -> None:\n  if not isinstance(v, ir.BlockArgument):\n    v.owner.attributes[name] = attr  # pyrefly: ignore[missing-attribute]\n    return","sourceCodeStart":1314,"sourceCodeEnd":1350,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/triton/lowering.py#L1314-L1350","documentation":"Raised by the Pallas→Triton lowering rule for jax.debug_print: after merging callback args, any remaining keyword arguments are rejected because the Triton backend only implements printing with positional values. The check mirrors a similar restriction against placeholder formatting. It is thrown at compile/lowering time, before the kernel runs on GPU.","triggerScenarios":"Calling pl.debug_print(fmt, **kwargs) (any keyword arguments) inside a Pallas kernel compiled with jax.experimental.pallas with triton as the target; also passing a pytree of arguments that merge_callback_args flattens into kwargs.","commonSituations":"Porting a Pallas/Mosaic kernel from the TPU backend (where debug_print accepts kwargs) to the Triton GPU backend; using named arguments out of habit from Python's print-style APIs; upgrading JAX versions where the restriction started being enforced.","solutions":["Pass all values to pl.debug_print positionally, e.g. pl.debug_print('x={} y={}', x, y) instead of pl.debug_print('x={x}', x=x)","Remove any placeholder-style format strings as well; pre-format into a simple '{}/{}/...' positional template","If you need named debugging, compute a tuple and print it positionally or fall back to jax.debug.print outside the kernel"],"exampleFix":"// before\npl.debug_print(\"a={a} b={b}\", a=block, b=acc)\n// after\npl.debug_print(\"a={} b={}\", block, acc)","handlingStrategy":"validation","validationCode":"def check_debug_print_args(fmt, args, kwargs):\n    if kwargs:\n        raise TypeError(\"pl.debug_print on Triton only supports positional args\")\n    if '{' in fmt and not re.fullmatch(r'(\\{\\})*', fmt.replace(' ', '')):\n        raise TypeError(\"use positional {} placeholders, not named placeholders\")\n    return True","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Always call pl.debug_print(fmt, *values) with positional args and '{}' placeholders","Wrap debug printing in a helper that enforces positional args so all call sites stay Triton-safe"],"tags":["jax","pallas","triton","debug-print","kwargs","gpu"],"backgroundTag":"backend-unsupported-argument","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}