jax-ml/jax · error · NotImplementedError

Not implemented: accuracy

Error message

Not implemented: accuracy

What it means

The Mosaic GPU lowering rule for lax.rsqrt accepts an accuracy hint argument but does not implement it; passing any non-None accuracy raises NotImplementedError('Not implemented: accuracy').

Source

Thrown at jax/_src/pallas/mosaic_gpu/lowering.py:2957

    return arith_dialect.mulf(x, x)
  raise NotImplementedError(f"Unsupported dtype {x_aval.dtype}")


@register_lowering_rule(lax.clz_p, mgpu.LoweringSemantics.Lane)
@register_lowering_rule(lax.clz_p, mgpu.LoweringSemantics.Warpgroup)
def _clz_lowering_rule(ctx: LoweringRuleContext, x):
  [x_aval] = ctx.avals_in
  if ctx.module_ctx.lowering_semantics == mgpu.LoweringSemantics.Lane:
    return _ensure_fa(x, x_aval.dtype)._pointwise(math_dialect.ctlz, restrict_bitwidth=False)
  x = _ensure_ir_value(x, x_aval.dtype)
  return math_dialect.ctlz(x)


@register_lowering_rule(lax.rsqrt_p, mgpu.LoweringSemantics.Lane)
@register_lowering_rule(lax.rsqrt_p, mgpu.LoweringSemantics.Warpgroup)
def _rsqrt_lowering_rule(ctx: LoweringRuleContext, x, accuracy):
  if accuracy is not None:
    raise NotImplementedError("Not implemented: accuracy")
  [x_aval] = ctx.avals_in
  if ctx.module_ctx.lowering_semantics == mgpu.LoweringSemantics.Lane:
    return _ensure_fa(x, x_aval.dtype).rsqrt(approx=ctx.module_ctx.approx_math)
  fastmath = (
      arith_dialect.FastMathFlags.afn if ctx.module_ctx.approx_math else None
  )
  return math_dialect.rsqrt(
      _ensure_ir_value(x, x_aval.dtype), fastmath=fastmath
  )


@register_lowering_rule(lax.tanh_p, mgpu.LoweringSemantics.Lane)
@register_lowering_rule(lax.tanh_p, mgpu.LoweringSemantics.Warpgroup)
def _tanh_lowering_rule(ctx: LoweringRuleContext, x, accuracy):
  if accuracy is not None:
    raise NotImplementedError("Not implemented: accuracy")
  [x_aval] = ctx.avals_in
  if ctx.module_ctx.lowering_semantics == mgpu.LoweringSemantics.Lane:

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Call lax.rsqrt without the accuracy argument and rely on ctx.module_ctx.approx_math to control fastmath behavior
2. Set the approx_math flag on the compilation context instead of per-op accuracy
3. Fall back to 1 / jnp.sqrt(x) if specific accuracy behavior is needed

Example fix

// before
r = lax.rsqrt(x, accuracy=1e-6)
// after
r = lax.rsqrt(x)  # fastmath controlled by module approx_math

# strip accuracy hints before kernel launch
def rsqrt_safe(x, accuracy=None):
  if accuracy is not None:
    raise ValueError('accuracy unsupported in Mosaic; use approx_math flag')
  return lax.rsqrt(x)

try:
  out = kernel(x)
except NotImplementedError as e:
  if 'accuracy' in str(e): call lax.rsqrt(x) without accuracy or use 1/jnp.sqrt

When it happens

Trigger: Calling lax.rsqrt(x, accuracy=...) (or an API that forwards an accuracy hint, e.g. some precision-controlled math frontends) inside a Mosaic GPU kernel.

Common situations: Using accuracy-controlled math APIs or porting code that requests approximate vs precise rsqrt semantics; the kernel compiler path rejects the accuracy contract entirely.

Related errors

• Not implemented: accuracy
• Unsupported dtype: {ref.dtype}
• Only SMEM and TMEM refs are supported.
• Unsupported transform: {type(transform)}
• Non-trivial indexing on WGMMAAbstractAccumulatorRef is not s

AI-assisted analysis of jax-ml/jax@1e1c6a8fc0 (2026-08-27). Data as JSON: /api/errors/260d87d37f42c2d4. Report an issue: GitHub.