Chapter 5: Interoperability and Custom Operations
Study notes for the course "Advanced JAX: Performance, Optimization, and Scale" · Chapter 5 Original course: https://apxml.com/courses/advanced-jax (this chapter:
chapter-5-jax-interoperability-custom-operations) Environment: Python 3.14 · JAX / jaxlib 0.12.0.1 (CPU only) · NumPy · SciPy · Apple clang++
How to read this chapter: This is the chapter where the API has changed the most. host_callback, to_dlpack, xla_client, register_translation and def_jvp/def_vjp no longer work in JAX 0.12. All the code in the main text uses the 0.12 style. Outdated code sits in ❌ cells, each followed directly by the ✅ recommended version, and every conclusion is checked with an assert.
⚠️ Semi-internal API warning: Sections 5.5–5.9 use jax.extend.core.Primitive, jax.interpreters.{mlir, ad, batching} and jax.core.ShapedArray. These are semi-internal APIs, and JAX doesn't promise they stay stable across versions. The import paths here are only guaranteed to work on 0.12.0.1. After upgrading JAX, run those cells first before reading on.
Learning goals¶
- Know which conversions between
jax.Arrayandnumpy.ndarraycopy, which share memory, and where the sync points are. - Use the DLPack protocol (
__dlpack__/np.from_dlpack/jax.dlpack.from_dlpack) for zero-copy exchange, and prove the memory really is shared by comparing pointers. - Learn the three replacements for the removed
host_callback:jax.debug.callback,io_callbackandpure_callback, and how each behaves under jit/vmap/grad. - Understand that a primitive is the smallest unit JAX transformations work on. Define one from scratch and add, step by step, its impl, abstract evaluation, lowering, JVP, transpose and batching rules, so it works with jit, grad and vmap.
- Use XLA FFI (
jax.ffi) to plug a real compiled C++ function into JAX, then add gradient and vmap support.
Sections¶
| Section | Topic | Original slug |
|---|---|---|
| 5.1 | Integrating JAX with NumPy | integrating-jax-numpy |
| 5.2 | Zero-copy data sharing with DLPack | zero-copy-dlpack |
| 5.3 | calling-external-code-host-callback | |
| 5.4 | Side-effect-free calls with jax.pure_callback | using-jax-pure-callback |
| 5.5 | Introduction to JAX primitives | introduction-jax-primitives |
| 5.6 | Defining custom primitives | defining-custom-primitives |
| 5.7 | Implementing abstract evaluation rules | implementing-abstract-evaluation |
| 5.8 | Implementing lowering rules for backends | implementing-lowering-rules |
| 5.9 | Differentiation rules for custom primitives (+ batching) | differentiation-rules-custom-primitives |
| 5.10 | Practice: integrating a C++ function (ctypes callback → XLA FFI) | practice-integrating-cpp-function |
After that come the hands-on exercises, the quiz, the version notes, the chapter summary and the cheat sheet.
Links to the Basics part¶
- Asynchronous dispatch and
block_until_ready: Basics 01 Introduction, Basics 02 JIT - jaxpr, tracing and tracers: Basics 02 JIT. Where each line of a jaxpr comes from becomes clear in 5.5–5.7, which cover primitives and abstract evaluation.
- JVP/VJP and
custom_jvp/custom_vjp: Basics 03 grad. Sections 5.4, 5.9 and 5.10 all rely on them to give a "black box" a gradient. - vmap's batch dimension: Basics 04 vmap. This chapter's
vmap_methodand batching rules are the machinery underneath vmap. - Main table of contents: Basics 00 Overview
⚠️ Environment note: the jaxlib 0.12.0.1 CPU backend has a known bug: when a jit closure captures a large constant matrix and does a reduce after a matmul (
W @ X), it silently returns wrong results (for a reproduction and workaround, see "Known issues in this environment" in Basics 00 Overview; for the root-cause analysis, see Advanced 02 §2.3). So in this chapter, weights and data are always passed in as function arguments, and key numbers are cross-checked against NumPy float64.
# Shared setup: imports, version check, C++ directory, timing and aligned-allocation helpers
import sys, time, ctypes, subprocess, warnings, shutil
from pathlib import Path
from functools import partial
import numpy as np
import jax
import jax.numpy as jnp
print("Python", sys.version.split()[0], "| JAX", jax.__version__, "| devices", jax.devices())
assert jax.__version__.startswith("0.12")
assert jax.default_backend() == "cpu", "The examples in this chapter are written for CPU (on Colab, use a CPU runtime)"
NB_DIR = Path.cwd() # directory of this notebook
CPP_DIR = NB_DIR / "cpp" # where C++ sources and .so files go
CPP_DIR.mkdir(exist_ok=True)
def short_path(s):
"""Shorten local absolute paths when printing (notebook dir → relative path, virtualenv → <venv>) so no personal directory shows up in the output"""
return str(s).replace(str(NB_DIR) + "/", "").replace(sys.prefix, "<venv>")
print("C++ directory:", short_path(CPP_DIR))
def bench(f, *args, n=20):
"""Average time (ms): warm up once, loop n times, then block_until_ready"""
jax.block_until_ready(f(*args))
t0 = time.perf_counter()
for _ in range(n):
out = f(*args)
jax.block_until_ready(out)
return (time.perf_counter() - t0) / n * 1e3
def aligned_zeros(n, dtype=np.float32, align=64):
"""Allocate a 1-D NumPy array whose start address is aligned to `align` bytes (required for zero-copy DLPack import into JAX)"""
itemsize = np.dtype(dtype).itemsize
buf = np.zeros(n * itemsize + align, dtype=np.uint8)
off = (-buf.ctypes.data) % align
out = buf[off:off + n * itemsize].view(dtype)
assert out.ctypes.data % align == 0
return out
def shares_memory(j, n):
"""Whether JAX array j and NumPy array n point to the same memory (compares start pointers)"""
return j.unsafe_buffer_pointer() == n.ctypes.data
5.1 Integrating JAX with NumPy¶
5.1 Integrating JAX with NumPy¶
jax.numpy deliberately mirrors the NumPy API, but the two kinds of arrays are fundamentally different, and converting between them has performance costs to watch for.
How the two arrays differ
numpy.ndarray: the data lives in host CPU memory, and operations run directly on the CPU (eagerly).jax.Array: the data is managed by the JAX runtime. It can live on the CPU, but its real strength is that it can live on accelerators like GPUs and TPUs. Transformations such asjit,gradandvmap, and compiled computations, all work on it.- Because the two live in different places and serve different purposes, conversions deserve some care.
NumPy → JAX
- The most direct way is
jnp.array(np_arr): JAX reads the data from host memory and puts it on the default device (a GPU/TPU if there is one, otherwise the CPU). - To say exactly which device it goes to, use
jax.device_put(np_arr, device=...). You need this for distributed computing, or when you have to guarantee data locality.
JAX → NumPy
np.array(jax_arr)pulls a result back to the host so you can save it, plot it or hand it to another library.- Watch the performance here: this conversion has two steps. ① Wait for the device to finish computing the array. ② Copy the data from device memory back to host memory. JAX uses asynchronous dispatch by default (Basics 01), so this step is a sync point: Python waits until the data is ready. Converting often inside a performance-sensitive loop can slow things down a lot.
The course's 5 best practices
- Convert as little as possible in inner loops, and never inside an
@jax.jitfunction. - Convert only at the boundaries of your computation: bring data in with
jnp.array/jax.device_putat the start, and take it out withnp.arrayat the end. - On multiple devices, use
jax.device_putto state placement explicitly. np.array(jax_array)quietly waits for the computation to finish. If you only want to sync and not copy data (for accurate timing, say), call.block_until_ready()directly.- Prefer
jax.numpyin JAX code. Passing NumPy arrays straight intojnpfunctions does work, but the implicit conversions can hide performance costs.
In short, treat conversions (especially JAX → NumPy) as potentially expensive, and decide deliberately where they go when you design a program.
# ✅ Recommended: NumPy → JAX (course code block 1; use the .device attribute for the device)
numpy_arr = np.array([1.0, 2.0, 3.0], dtype=np.float32)
print("Original NumPy type:", type(numpy_arr), "(implicitly on the host CPU)")
jax_arr = jnp.array(numpy_arr)
print("Converted type:", type(jax_arr), "| isinstance(jax.Array):", isinstance(jax_arr, jax.Array))
print("Device:", jax_arr.device)
assert isinstance(jax_arr, jax.Array) and jax_arr.dtype == jnp.float32
assert jax_arr.device == jax.devices()[0]
# ❌ Old style (outdated/wrong): jax_arr.device() — in 0.12, .device is an attribute (a Device object) and can't be called
try:
print(jax_arr.device())
except TypeError as e:
print("TypeError:", e)
# ✅ Recommended: the .device attribute (single-device arrays); .devices() returns the set of devices (works for sharded arrays too)
print(jax_arr.device, jax_arr.devices())
assert jax_arr.devices() == {jax_arr.device}
# Course code block 2: explicit placement with jax.device_put
available_devices = jax.devices()
print("Available devices:", available_devices)
target_device = available_devices[1] if len(available_devices) > 1 else available_devices[0]
jax_arr_explicit = jax.device_put(numpy_arr, device=target_device)
print("Device after explicit placement:", jax_arr_explicit.device)
assert jax_arr_explicit.device == target_device
# ❌ Old style (outdated/wrong): jax.random.PRNGKey(0) returns a uint32[2] "legacy raw key"; 0.12 recommends typed keys
try:
old_key = jax.random.PRNGKey(0)
print("Legacy key:", old_key, "| dtype:", old_key.dtype, "| shape:", old_key.shape)
except Exception as e:
print(type(e).__name__, e)
# ✅ Recommended: jax.random.key creates a typed key; JAX → NumPy (course code block 3)
key = jax.random.key(0)
print("Typed key:", key, "| dtype:", key.dtype)
assert jax.dtypes.issubdtype(key.dtype, jax.dtypes.prng_key)
jax_arr = jax.random.normal(key, (3,)) * 2.0 + 1.0
print("JAX array:", jax_arr, "| device:", jax_arr.device)
numpy_result = np.array(jax_arr) # sync point: waits for the computation, then copies back to the host
print("NumPy type:", type(numpy_result), "| value:", numpy_result)
assert isinstance(numpy_result, np.ndarray)
np.testing.assert_array_equal(numpy_result, jax_arr)
⚠️ Note:
jnp.arrayandjax.device_puthave different memory semantics. The difference is more than "one uses the default device and the other takes an explicit one". On the JAX 0.12 CPU backend:
jnp.array(np_arr)/jnp.asarray(np_arr): always copy.jax.device_put(np_arr): as long as the source array is 64-byte aligned, it aliases the NumPy memory directly, and passingmay_alias=Falsedoesn't stop it. On this macOS machine, NumPy arrays of a few KB or more are 64-byte aligned; very small arrays may not be, depending on the platform's memory allocator. If you then modify the NumPy source array, this "immutable" JAX array changes along with it.- Going the other way,
np.asarray(jax_arr)is a zero-copy read-only view; onlynp.array(jax_arr)copies.
# Try it: which conversions actually copy (decided by comparing pointers)
big = aligned_zeros(1 << 20); big[:] = 1.0 # 4 MB, 64-byte aligned
a_copy = jnp.array(big)
a_asarray = jnp.asarray(big)
a_put = jax.device_put(big)
a_put_noalias = jax.device_put(big, may_alias=False)
for name, arr in [("jnp.array", a_copy), ("jnp.asarray", a_asarray),
("jax.device_put", a_put), ("device_put(may_alias=False)", a_put_noalias)]:
print(f"{name:28s} shares memory with NumPy: {shares_memory(arr, big)}")
big[0] = 42.0 # modify the NumPy source array
print("After modifying the source: jnp.array[0] =", float(a_copy[0]), "| device_put[0] =", float(a_put[0]))
assert not shares_memory(a_copy, big) and float(a_copy[0]) == 1.0
assert not shares_memory(a_asarray, big)
assert shares_memory(a_put, big) and float(a_put[0]) == 42.0 # ⚠️ the "immutable" JAX array changed
assert shares_memory(a_put_noalias, big) # measured on 0.12.0.1: may_alias=False aliases too
# JAX → NumPy
y = jax.jit(lambda v: v * 2)(jnp.arange(1024.0))
n_view, n_copy = np.asarray(y), np.array(y)
print("np.asarray zero-copy:", n_view.ctypes.data == y.unsafe_buffer_pointer(), "| writeable:", n_view.flags.writeable)
print("np.array zero-copy:", n_copy.ctypes.data == y.unsafe_buffer_pointer(), "| writeable:", n_copy.flags.writeable)
assert n_view.ctypes.data == y.unsafe_buffer_pointer() and not n_view.flags.writeable
assert n_copy.ctypes.data != y.unsafe_buffer_pointer() and n_copy.flags.writeable
# The cost of sync points: convert back to NumPy every step vs. only once at the end (timings are printed, not asserted)
@jax.jit
def step(x):
return jnp.sin(x) * 0.999 + 0.001
x0 = jnp.ones((512, 512))
def loop_convert_each_step(x, n=100):
hist = []
for _ in range(n):
x = step(x)
hist.append(np.array(x)[0, 0]) # sync and copy every step
return x
def loop_convert_at_end(x, n=100):
for _ in range(n):
x = step(x) # queued asynchronously, doesn't block
return x
t_each = bench(loop_convert_each_step, x0, n=3)
t_end = bench(loop_convert_at_end, x0, n=3)
print(f"convert every step: {t_each:.2f} ms | convert only at the end: {t_end:.2f} ms | ratio {t_each / t_end:.2f}x")
print("Reading: " + ("syncing every step is slower, as expected." if t_each > t_end * 1.1 else
"CPU only here: device memory is host memory, so \"copying back\" is nearly free and async queuing gains little. The gap is small and may even flip. "
"The cost of sync points mostly shows up on GPU/TPU, where every sync waits for the device and goes over PCIe."))
np.testing.assert_allclose(loop_convert_each_step(x0), loop_convert_at_end(x0))
🔍 Extra¶
🟢 Must-read
- You can't apply NumPy functions to tracers inside jit:
np.sin(tracer)raisesTracerArrayConversionError. NumPy needs concrete values, but during tracing there are only abstract values (see Basics 02 JIT). If you really need NumPy/SciPy inside jit, use the callbacks covered in 5.3 and 5.4. - Zero-copy views cut both ways:
np.asarray(jax_arr)doesn't copy, but it's read-only.jax.device_put(np_arr)may alias on CPU, so if you still plan to modify that NumPy array in place afterwards, make anp_arr.copy()first, or usejnp.arrayinstead. jax.device_get(tree)pulls a whole pytree back to the host in one call, which is easier than callingnp.arrayon each leaf.
# Check: calling a NumPy function on a tracer inside jit
@jax.jit
def bad(x):
return np.sin(x) # during tracing x is a tracer, and NumPy can't turn it into a concrete array
try:
bad(jnp.ones(3))
err = None
except Exception as e:
err = e
print(type(e).__name__, ":", str(e).splitlines()[0][:120])
assert isinstance(err, jax.errors.TracerArrayConversionError)
host = jax.device_get({"w": jnp.ones(2), "b": jnp.zeros(1)})
assert all(isinstance(v, np.ndarray) for v in host.values())
5.2 Zero-copy data sharing with DLPack¶
5.2 Zero-copy data sharing with DLPack¶
The problem: the usual way to exchange data between JAX and NumPy, PyTorch, TensorFlow or CuPy is to copy it, for example JAX → host NumPy → PyTorch → back to the GPU. For large arrays on an accelerator, all this shuffling and duplicated memory is expensive. Even when JAX and PyTorch sit on the same GPU, copying several GB between their two memory pools wastes time and eats precious memory bandwidth.
The DLPack standard: a language-independent specification for describing "a tensor in memory". When the data is on the same device, libraries that support it can exchange data with zero copies; all that actually gets passed is a pointer plus metadata.
- The producer creates a capsule (a
PyCapsulein Python) that holds: ① a pointer to the underlying buffer; ② metadata: device type and ID (CPU/CUDA/ROCm…), dtype, number of dimensionsndim, shape, strides and byte offset. - The consumer reads the pointer and metadata and wraps the existing memory directly in its own tensor object, without copying the data.
The course's 5 caveats
- Device compatibility: zero-copy only works when both sides use the same physical device. DLPack can't share directly between CPU and GPU; the data has to be moved to the target device first.
- Synchronization: JAX executes asynchronously, so call
block_until_ready()before exporting. Before importing from another framework (say, PyTorch on CUDA), also make sure its compute stream has synced, or you might read data that hasn't finished being written. - Lifetime and ownership: the memory usually stays owned by the original framework. The original tensor must stay alive while the imported object is in use, or you get a dangling pointer.
- Mutability: JAX arrays are immutable. If the consumer modifies the underlying buffer in place, it breaks that assumption.
- Stream ordering (advanced): on a GPU, if the two sides use different CUDA streams, calling
block_until_ready()ortorch.cuda.synchronize()may not be enough to make the data visible to the other side; you need stream events to coordinate.
⚠️ Note: a common pattern in older material is to export a capsule with
jax.dlpack.to_dlpack(x)and then pass it totorch.from_dlpackorjax.dlpack.from_dlpack, with the warning that "a capsule can only be consumed once". In JAX 0.12,jax.dlpack.to_dlpackhas been removed; the module only hasfrom_dlpackandis_supported_dtype.jax.dlpack.from_dlpackalso no longer accepts raw capsules; passing one raisesTypeError.The current way is the DLPack protocol: hand the array object directly to the consumer, for example
np.from_dlpack(x)ortorch.from_dlpack(x), and the consumer callsx.__dlpack__()andx.__dlpack_device__()itself. You never touch the capsule, so the "use it only once" issue goes away. Also,x.device_buffer.unsafe_buffer_pointer()has been removed; usex.unsafe_buffer_pointer()orx.addressable_data(0).unsafe_buffer_pointer()instead.
# ❌ Old style (outdated/wrong): to_dlpack / device_buffer / from_dlpack(capsule) / random.normal(device=) all fail in 0.12
x_src = jax.random.normal(jax.random.key(0), (1024, 1024))
attempts = {
"jax.dlpack.to_dlpack(x)": lambda: jax.dlpack.to_dlpack(x_src),
"x.device_buffer.unsafe_buffer_pointer()": lambda: x_src.device_buffer.unsafe_buffer_pointer(),
"jax.dlpack.from_dlpack(x.__dlpack__())": lambda: jax.dlpack.from_dlpack(x_src.__dlpack__()),
"jax.random.normal(key, shape, device=...)": lambda: jax.random.normal(jax.random.key(0), (2,), device=jax.devices()[0]),
}
for desc, fn in attempts.items():
try:
fn(); print(f"{desc:45s} → still works")
except Exception as e:
print(f"{desc:45s} → {type(e).__name__}: {str(e).splitlines()[0][:90]}")
print("Public members of jax.dlpack:", [a for a in dir(jax.dlpack) if not a.startswith("_")])
print("(torch / cupy aren't installed here, so those two examples can't run in this environment)")
# ✅ Recommended: the DLPack protocol — zero-copy JAX → NumPy
x = jax.random.normal(jax.random.key(0), (1024, 1024))
x.block_until_ready() # sync before exporting (caveat 2)
print("__dlpack_device__:", x.__dlpack_device__())
print("__dlpack__() returns:", type(x.__dlpack__()).__name__)
n = np.from_dlpack(x) # NumPy calls x.__dlpack__() itself
print("JAX pointer :", x.unsafe_buffer_pointer())
print("NumPy pointer:", n.ctypes.data, "| writeable:", n.flags.writeable)
assert not hasattr(jax.dlpack, "to_dlpack")
assert n.ctypes.data == x.unsafe_buffer_pointer() == x.addressable_data(0).unsafe_buffer_pointer()
assert not n.flags.writeable # the export is a read-only view, which protects JAX's immutability
np.testing.assert_array_equal(n, x)
# ✅ Recommended: zero-copy NumPy → JAX (jax.dlpack.from_dlpack / jnp.from_dlpack)
src = aligned_zeros(1 << 20); src[:] = np.arange(src.size, dtype=np.float32)
j1 = jax.dlpack.from_dlpack(src)
j2 = jnp.from_dlpack(src) # equivalent entry point
print("from_dlpack shares memory:", shares_memory(j1, src), "| jnp.from_dlpack shares memory:", shares_memory(j2, src))
assert shares_memory(j1, src) and shares_memory(j2, src)
# Alignment requirement: if the start address isn't 64-byte aligned, the default (copy=None) silently copies; copy=False raises
unaligned = aligned_zeros(1025)[1:] # offset by 4 bytes
assert unaligned.ctypes.data % 64 != 0
j_u = jax.dlpack.from_dlpack(unaligned)
print("unaligned array, from_dlpack shares memory:", shares_memory(j_u, unaligned))
assert not shares_memory(j_u, unaligned)
try:
jax.dlpack.from_dlpack(unaligned, copy=False)
raised = False
except ValueError as e:
raised = True
print("copy=False →", str(e)[:110])
assert raised
⚠️ Note: a JAX array imported with zero copies isn't truly "unchangeable". You can't change it through JAX's API, but it shares the same memory as the source NumPy array: modify the source array directly and the JAX array's values change too. That breaks JAX's functional assumptions, for example constants cached by jit, or values read by async computations already in the queue. So after a zero-copy import, treat the source array as "frozen" and leave it alone.
# Check: after a zero-copy import, modifying the source array changes the JAX array; lifetime is guaranteed by the consumer holding a reference
src = aligned_zeros(1024); src[:] = 1.0
j = jax.dlpack.from_dlpack(src)
before = float(j[0])
src[0] = 99.0
print("before: j[0] =", before, "| after modifying the source: j[0] =", float(j[0]))
assert before == 1.0 and float(j[0]) == 99.0
# Lifetime: after deleting the JAX array, the NumPy view is still valid (DLPack's deleter mechanism makes the consumer hold a reference to the producer)
x = jnp.full((1000,), 3.0)
view = np.from_dlpack(x)
del x
import gc; gc.collect()
print("After deleting the original JAX array, NumPy view sum =", view.sum())
assert view.sum() == 3000.0
# Timing: zero-copy vs. copy (64 MB; only the comparison is printed)
big = aligned_zeros(16 << 20); big[:] = 1.0
t_copy = bench(lambda a: jnp.array(a), big, n=10)
t_dl = bench(lambda a: jax.dlpack.from_dlpack(a), big, n=10)
print(f"jnp.array (copy): {t_copy:.3f} ms | from_dlpack (zero-copy): {t_dl:.4f} ms | {t_copy / t_dl:.0f}x faster")
🔍 Extra¶
⚪ Optional: when you need to exchange data with frameworks like PyTorch or CuPy
- The modern way to convert to and from PyTorch (torch isn't installed here, so this wasn't actually run):
torch.from_dlpack(jax_arr)andjax.dlpack.from_dlpack(torch_tensor)both take the object directly, no intermediate capsule needed. On GPU, the consumer negotiates stream synchronization through__dlpack__(stream=...). - Supported dtypes: check with
jax.dlpack.is_supported_dtype(dtype). - Across devices:
jax.dlpack.from_dlpack(obj, device=..., copy=...)lets you pick the target device. If the target device differs from the source device, it always copies, and passingcopy=Falsethen raises. Put the other way round:copy=Falseguarantees you really get zero-copy; without it, a zero-copy import may quietly turn into a copy.
5.3 Calling external CPU/GPU code: from host_callback to three kinds of callbacks¶
5.3 Calling external CPU/GPU code with jax.experimental.host_callback¶
XLA compilation is what makes JAX fast. But sometimes we need to step out of the compiled world while a computation is running and run some arbitrary Python on the host CPU: for debugging, logging intermediate values, talking to external hardware, or calling a library that has no JAX implementation. The course covers jax.experimental.host_callback, which mainly offers two functions:
id_tap(fn, arg, tap_with_transform=...): acts as an identity function in the computation graph (returnsargunchanged) while triggeringfnon the host, for side effects like printing and logging.id_printis a simplified version just for printing.call(fn, args, result_shape_dtypes=...): calls a host function and sends its return value back to the device, so it rejoins JAX's data flow. The output shape and dtype must be known at tracing time, soresult_shape_dtypesis required.
How it works (at execution time, not tracing time): ① when execution reaches the callback instruction, the device pauses; ② the arguments are copied from device to host; ③ the Python function runs on the host; ④ (only for call) the return value is copied back to the device; ⑤ the device continues. This round trip brings a sync point and transfer overhead.
The course's caveats: ① high performance overhead; ② to XLA it's an opaque optimization barrier, so ops on either side of the callback can't be fused; ③ callbacks are dispatched asynchronously too, so prints may only appear once the result is needed; ④ side effects are allowed, so the function isn't pure; ⑤ not differentiable by default; ⑥ interaction with transformations: it works under jit, behaves in complicated ways under vmap, and runs once per host under pmap.
When to use it: as a last resort, or for tasks where performance doesn't matter, such as debugging, logging, hooking up legacy code or talking to hardware. Alternatives: rewrite with jax.numpy/jax.lax; use jax.pure_callback for purely functional external calls (5.4); use custom primitives or FFI for high-performance C++/CUDA (5.6–5.10).
⚠️ Note:
jax.experimental.host_callback(id_tap,id_print,call) has been removed in JAX 0.12.importlib.util.find_specreturnsNone, and importing it raisesImportError. Its features were split across three official APIs: | Old API | 0.12 replacement | Use case | |---|---|---| |id_tap/id_print|jax.debug.callback/jax.debug.print| Debug printing, no return value | |call(with side effects) |jax.experimental.io_callback(..., ordered=True/False)| Has side effects, needs a return value or guaranteed ordering | |call(pure function) |jax.pure_callback(..., vmap_method=...)| Pure computation, can be vmapped, can be removed by DCE |
tap_with_transformhas no replacement. Also, it's easy to assume the host function receives NumPy arrays, but on the 0.12 CPU backend all three callbacks actually receivejax.Array(ArrayImpl). So inside a callback, donp.asarray(arg)first before handing it to NumPy/SciPy.
# ❌ Old style (outdated/wrong): from jax.experimental import host_callback — the module was removed in 0.12
import importlib.util
print("find_spec:", importlib.util.find_spec("jax.experimental.host_callback"))
try:
from jax.experimental import host_callback
host_callback.id_tap(print, jnp.ones(2))
except ImportError as e:
print("ImportError:", short_path(e))
assert importlib.util.find_spec("jax.experimental.host_callback") is None
# ✅ Recommended: id_tap → jax.debug.callback (course code block 1)
records = []
def log_intermediate_data(arg):
records.append(type(arg))
arg = np.asarray(arg) # the callback receives a jax.Array; convert to NumPy first
print(f"[host callback] shape={arg.shape} first 5={np.round(arg.ravel()[:5], 4)}")
@jax.jit
def process_data(x):
y = jnp.sin(x) * 2.0
jax.debug.callback(log_intermediate_data, y) # no return value, doesn't affect the data flow
z = jnp.mean(jnp.cos(y))
return z
input_data = jax.random.uniform(jax.random.key(42), (8, 8))
result = process_data(input_data)
result.block_until_ready()
jax.effects_barrier() # wait for all callbacks so prints don't leak into other cells
print("Final result:", result)
assert len(records) == 1 and issubclass(records[0], jax.Array)
ref = np.mean(np.cos(np.sin(np.asarray(input_data, np.float64)) * 2.0))
np.testing.assert_allclose(result, ref, rtol=1e-5)
# ✅ Recommended: id_print → jax.debug.print (format string, formatted lazily)
@jax.jit
def simple_print_example(x):
y = x + 5.0
jax.debug.print("intermediate y = {y}", y=y)
return y * 2.0
output = simple_print_example(jnp.arange(3.0))
output.block_until_ready(); jax.effects_barrier()
np.testing.assert_array_equal(output, [10.0, 12.0, 14.0])
# ✅ Recommended: host_callback.call (pure computation) → jax.pure_callback (course code block 2)
def external_cpu_calculation(data, parameter):
data = np.asarray(data) # host side: ordinary NumPy computation
return (np.tanh(data) + np.asarray(parameter)).astype(data.dtype)
@jax.jit
def jax_workflow(x, ext_param):
intermediate = jnp.log1p(jnp.abs(x))
external_result = jax.pure_callback(
external_cpu_calculation,
jax.ShapeDtypeStruct(intermediate.shape, intermediate.dtype), # output shape and dtype
intermediate, ext_param)
return external_result / (1.0 + jnp.mean(intermediate))
input_array = jnp.linspace(-2.0, 2.0, 5, dtype=jnp.float32)
final_result = jax_workflow(input_array, 0.5)
print("Final result:", final_result)
xi = np.log1p(np.abs(np.linspace(-2, 2, 5)))
np.testing.assert_allclose(final_result, (np.tanh(xi) + 0.5) / (1 + xi.mean()), rtol=1e-5)
# ✅ Recommended: calls with side effects → io_callback; ordered=True keeps execution order the same as program order
from jax.experimental import io_callback
order = []
def write_log(i):
order.append(int(i)) # side effect: writes external state
def body(i, carry):
io_callback(write_log, None, i, ordered=True) # when the callback returns None, pass None for result_shape_dtypes
return carry + i
total = jax.jit(lambda: jax.lax.fori_loop(0, 5, body, 0))()
jax.effects_barrier()
print("Write order:", order, "| total =", total)
assert order == [0, 1, 2, 3, 4] and int(total) == 10
How the three callbacks behave under jit / vmap / grad / DCE (each case tested)¶
The next cell tries each callback under each transformation one by one, produces the table below, and pins the conclusions down with assert. The "when the output is unused" column is about DCE (dead code elimination): if a callback's output isn't used, does jit delete the call as dead code?
| Callback | jit | vmap | grad | When the output is unused |
|---|---|---|---|---|
jax.debug.callback | ✅ | ✅ (called per example, B times) | ✅ (sees only primal values, contributes nothing to the gradient) | Kept (has side effects) |
io_callback(ordered=False) | ✅ | ✅ (called per example, B times) | ❌ ValueError: JVP not supported | Kept |
io_callback(ordered=True) | ✅ | ❌ ValueError: can't vmap an ordered callback | ❌ ValueError | Kept |
pure_callback (no vmap_method) | ✅ | ❌ NotImplementedError | ❌ ValueError: suggests custom_jvp | Deleted, never runs |
pure_callback(vmap_method="expand_dims") | ✅ | ✅ (one call that receives the whole batch) | ❌ ValueError | Deleted |
# Test each entry of the table above
def run(fn):
try:
jax.block_until_ready(fn()); jax.effects_barrier()
return "OK"
except Exception as e:
return type(e).__name__
calls = []
def host_log(a): calls.append(np.shape(a))
def host_double(a):
calls.append(np.shape(a)); return np.asarray(a) * 2
sds = lambda v: jax.ShapeDtypeStruct(v.shape, v.dtype)
variants = {
"debug.callback": lambda v: (jax.debug.callback(host_log, v), v * 2)[1],
"io_callback": lambda v: io_callback(host_double, sds(v), v),
"io_callback(ordered)": lambda v: io_callback(host_double, sds(v), v, ordered=True),
"pure_callback(default)": lambda v: jax.pure_callback(host_double, sds(v), v),
"pure_callback(expand_dims)": lambda v: jax.pure_callback(host_double, sds(v), v, vmap_method="expand_dims"),
}
x3, xb = jnp.arange(3.0), jnp.ones((4, 3))
table = {}
for name, f in variants.items():
row = {}
row["jit"] = run(lambda: jax.jit(f)(x3))
calls.clear(); row["vmap"] = run(lambda: jax.vmap(f)(xb)); row["vmap calls"] = len(calls)
row["grad"] = run(lambda: jax.grad(lambda v: f(v).sum())(x3))
calls.clear(); jax.jit(lambda v: (f(v), v + 1)[1])(x3).block_until_ready(); jax.effects_barrier()
row["runs when unused"] = len(calls)
table[name] = row
print(f"{'callback':28s}{'jit':6s}{'vmap':22s}{'vmap calls ':10s}{'grad':12s}{'runs when unused'}")
for name, r in table.items():
print(f"{name:28s}{r['jit']:6s}{r['vmap']:22s}{r['vmap calls']!s:14s}{r['grad']:12s}{r['runs when unused']}")
assert all(r["jit"] == "OK" for r in table.values())
assert table["debug.callback"]["vmap"] == "OK" and table["debug.callback"]["vmap calls"] == 4
assert table["debug.callback"]["grad"] == "OK"
assert table["io_callback"]["vmap"] == "OK" and table["io_callback"]["grad"] == "ValueError"
assert table["io_callback(ordered)"]["vmap"] == "ValueError"
assert table["pure_callback(default)"]["vmap"] == "NotImplementedError"
assert table["pure_callback(default)"]["grad"] == "ValueError"
assert table["pure_callback(expand_dims)"]["vmap"] == "OK" and table["pure_callback(expand_dims)"]["vmap calls"] == 1
assert table["debug.callback"]["runs when unused"] == 1 and table["io_callback"]["runs when unused"] == 1
assert table["pure_callback(default)"]["runs when unused"] == 0 # the pure callback is removed as dead code
🔍 Extra¶
🟢 Must-read
- How to choose: to just peek at an intermediate value, use
jax.debug.printorjax.debug.callback. To write files, send metrics, or rely on execution order, useio_callback(ordered=True). For pure computation that needs a return value, usepure_callback.pure_callbackis the only callback that JAX is allowed to delete, merge or run more than once, because JAX trusts that it's pure. If it has side effects, the results are unpredictable (see the DCE check in 5.4). - Under grad,
debug.callbackonly sees primal values: it doesn't take part in differentiation and sees neither tangents nor cotangents. To print a gradient, calljax.debug.printon the result ofjax.grad. - Asynchrony: callbacks run asynchronously with device computation. In a notebook, printed output can "leak" into later cells;
jax.effects_barrier()waits until every callback has run. - Making a callback differentiable: none of the three callbacks supports autodiff on its own. Combine them with
jax.custom_jvporjax.custom_vjp; see 5.4.
5.4 Side-effect-free calls with jax.pure_callback¶
5.4 Side-effect-free calls with jax.pure_callback¶
When you want to plug an external Python function that does pure computation into a graph that needs jit, vmap or grad, use jax.pure_callback.
The purity contract: the wrapped function must be functionally pure. That is, it's deterministic (the same inputs always give the same outputs), and it has no side effects (it doesn't modify external state or do I/O, and its result depends only on the arguments passed in). JAX fully trusts you to keep this contract and never checks it. If you break it, the results may be wrong, usually without any clear error.
How it works: a call takes three things: the callback function, result_shape_dtypes describing the output structure, and the actual arguments. During tracing the callback is not executed; JAX just puts a placeholder pure_callback equation in the jaxpr and treats the callback as a black box. Only at run time is it called with concrete values.
How to write result_shape_dtypes: for a single output, use jax.ShapeDtypeStruct(shape, dtype). For multiple outputs, use a tuple with one ShapeDtypeStruct per element. Any pytree works, as long as its structure matches the callback's return value.
Use cases: hooking up external libraries such as SciPy; implementing complex logic in plain Python; filling in temporarily when JAX lacks some operation.
The course's caveats: ① you are responsible for purity (this one matters most); ② switching back and forth between accelerator and host and moving data has overhead, so native XLA ops are usually much faster; ③ compatibility with transformations; ④ functions that are closures or depend on external modules may be hard to serialize.
⚠️ Note: it's easy to assume pure_callback just works with jit, vmap and grad, for instance thinking "the callback computes sin(x)*y, which is differentiable, so
jax.gradshould work directly". In JAX 0.12 that isn't the case:
- vmap:
vmap_methoddefaults toNone, and then vmap raisesNotImplementedErrorright away, so you must set it explicitly. The valid values in 0.12 are"sequential","sequential_unrolled","expand_dims","broadcast_all"and"legacy_vectorized".- grad: whether or not the math inside the callback is differentiable, you get
ValueError: Pure callbacks do not support JVP. JAX can't see what NumPy computes inside, so you must provide the derivative by hand withjax.custom_jvporjax.custom_vjp.- The old argument
vectorized=Trueis no longer a pure_callback parameter in 0.12. It gets forwarded unchanged to the callback as part of**kwargs, and you getTypeError: ... unexpected keyword argument 'vectorized'at run time (see 5.10).
# ✅ Recommended: the course's main example (jit part), with the key created by jax.random.key
def external_pure_python_computation(x, y):
"""Plain Python/NumPy function: deterministic, no side effects"""
x = np.asarray(x) # the callback receives a jax.Array; convert to NumPy
return (np.sin(x) * y + 1.0).astype(x.dtype)
def jax_function_using_callback(a, b):
output_shape_dtype = jax.ShapeDtypeStruct(a.shape, a.dtype)
result = jax.pure_callback(
external_pure_python_computation, output_shape_dtype, a, b,
vmap_method="expand_dims") # declare the vmap semantics explicitly (see the comparison below)
return result + a
key = jax.random.key(0)
input_array = jax.random.normal(key, (3, 3))
input_scalar = 2.0
output = jax.jit(jax_function_using_callback)(input_array, input_scalar)
expected = np.sin(np.asarray(input_array, np.float64)) * input_scalar + 1.0 + np.asarray(input_array, np.float64)
print("Output:\n", output)
np.testing.assert_allclose(output, expected, rtol=1e-5)
# vmap: with expand_dims the callback is called only once and receives the whole (10, 3, 3) batch
batched_input = jax.random.normal(key, (10, 3, 3))
batched_output = jax.vmap(jax_function_using_callback, in_axes=(0, None))(batched_input, input_scalar)
print("Batched output shape:", batched_output.shape)
assert batched_output.shape == (10, 3, 3)
np.testing.assert_allclose(batched_output, np.sin(np.asarray(batched_input)) * 2 + 1 + np.asarray(batched_input), rtol=1e-5)
# ❌ Old style (outdated/wrong): vmap without vmap_method; grad straight through pure_callback
def course_version(a, b):
out = jax.pure_callback(external_pure_python_computation, jax.ShapeDtypeStruct(a.shape, a.dtype), a, b)
return out + a
try:
jax.vmap(course_version, in_axes=(0, None))(batched_input, input_scalar)
except NotImplementedError as e:
print("vmap → NotImplementedError:", str(e)[:150])
try:
jax.grad(lambda x: jnp.sum(course_version(x, input_scalar)))(input_array)
except ValueError as e:
print("grad → ValueError:", str(e)[:150])
# ✅ Recommended: give pure_callback a derivative with custom_jvp (d/dx [sin(x)·y + 1] = cos(x)·y)
@jax.custom_jvp
def ext_fn(a, b):
return jax.pure_callback(external_pure_python_computation, jax.ShapeDtypeStruct(a.shape, a.dtype),
a, b, vmap_method="expand_dims")
@ext_fn.defjvp
def ext_fn_jvp(primals, tangents):
a, b = primals
a_dot, b_dot = tangents
y = ext_fn(a, b)
# write the derivative in jnp, so it is itself differentiable, jittable and vmappable
return y, jnp.cos(a) * b * a_dot + jnp.sin(a) * b_dot
grad_fn = jax.jit(jax.grad(lambda x: jnp.sum(ext_fn(x, input_scalar) + x)))
gradients = grad_fn(input_array)
print("Gradients:\n", gradients)
np.testing.assert_allclose(gradients, np.cos(np.asarray(input_array)) * 2.0 + 1.0, rtol=1e-5)
# derivative with respect to the scalar argument b
gb = jax.grad(lambda b: jnp.sum(ext_fn(input_array, b)))(2.0)
np.testing.assert_allclose(gb, np.sin(np.asarray(input_array)).sum(), rtol=1e-5)
# vmap ∘ grad composes too
vg = jax.vmap(jax.grad(lambda x: jnp.sum(ext_fn(x, 2.0))))(batched_input)
np.testing.assert_allclose(vg, np.cos(np.asarray(batched_input)) * 2.0, rtol=1e-5)
# The course's multi-output example: result_shape_dtypes described by a tuple (pytree)
def multi_output_pure_function(x):
x = np.asarray(x)
return np.sum(x).astype(x.dtype), np.mean(x).astype(x.dtype)
def jax_multi_output_callback(arr):
output_structure = (jax.ShapeDtypeStruct((), arr.dtype), # scalar shape is ()
jax.ShapeDtypeStruct((), arr.dtype))
sum_val, mean_val = jax.pure_callback(multi_output_pure_function, output_structure, arr)
return sum_val * 2, mean_val * 3
res1, res2 = jax.jit(jax_multi_output_callback)(jnp.arange(5.0))
print(f"Multiple outputs: {res1=}, {res2=}")
assert float(res1) == 20.0 and float(res2) == 6.0
Values of vmap_method (tested on 0.12)¶
Let the batch size be $B$ and the input shape the callback normally handles be $s$:
vmap_method | How many times the callback runs | Shape the callback receives | When to use |
|---|---|---|---|
"sequential" | $B$ times (lax.scan loops over examples) | $s$ | The callback can only handle a single example |
"sequential_unrolled" | $B$ times (fully unrolled scan, unroll=B) | $s$ | Same as above, when B is small |
"expand_dims" | Once | Batched arguments are $(B, *s)$; unbatched arguments get a leading dimension of size 1 | The callback supports NumPy broadcasting (the most common choice) |
"broadcast_all" | Once | All arguments are broadcast to $(B, *s)$ | The callback needs every input to have the same batch dimension |
"legacy_vectorized" | Once | Batched arguments are $(B, *s)$, unbatched arguments keep their original shape (the old vectorized=True behavior) | Only for migrating old code |
Whichever you pick, the batch dimension is always first, and the first dimension of the callback's return value must also be $B$.
# Try each vmap_method: number of calls and the shapes the callback receives
seen = []
def h2(a, b):
seen.append((np.shape(a), np.shape(b)))
return (np.asarray(a) * np.asarray(b)).astype(np.float32)
def call_with(method):
def f(a):
return jax.pure_callback(h2, jax.ShapeDtypeStruct(a.shape, a.dtype), a, jnp.ones(3), vmap_method=method)
return f
results = {}
for m in ["sequential", "sequential_unrolled", "expand_dims", "broadcast_all", "legacy_vectorized"]:
seen.clear()
out = jax.vmap(call_with(m))(jnp.ones((4, 3)))
results[m] = list(seen)
print(f"{m:20s} called {len(seen)} time(s), (a, b) shapes received = {seen[0]} output {out.shape}")
assert len(results["sequential"]) == 4 and results["sequential"][0] == ((3,), (3,))
assert results["expand_dims"] == [((4, 3), (1, 3))] # the unbatched b just gets one extra size-1 dim
assert results["broadcast_all"] == [((4, 3), (4, 3))] # b is broadcast to the batch size
assert results["legacy_vectorized"] == [((4, 3), (3,))] # b keeps its original shape
try:
jax.vmap(call_with("bogus"))(jnp.ones((4, 3)))
except ValueError as e:
print("Invalid value →", str(e)[:120])
🔍 Extra¶
🟢 Must-read
- DCE and purity: when a pure_callback's output isn't used, jit deletes the whole call, and the callback never runs at all (see the table in 5.3). JAX may also call it more than once inside a transformation; for example, vmap with
sequentialcalls it B times. So never log, count or do I/O inside a pure_callback. - Shape and dtype must match: if the returned shape is wrong, you get
JaxRuntimeError: CpuCallback errorat run time. If the dtype is wrong, sometimes it raises (for example, declaring int32 but returning float), and sometimes it's quietly converted (with x64 off, float64 is converted to float32). The safe habit is to end the callback with an explicit.astype(<declared dtype>). - Higher-order derivatives: if the derivative inside
custom_jvpis also computed with a pure_callback, first-order grad works, butjax.hessianruns into "JVP not supported" again. To support higher-order derivatives, the derivative function itself must be differentiable, for example written in jnp, or wrapped in another custom_jvp. The cell below verifies this. - Performance: every call crosses between XLA and Python once, which has a fixed overhead and also blocks fusion. The timing comparison below shows this.
# Check: when the derivative also uses pure_callback, hessian fails; wrapping the derivative in custom_jvp too fixes it
np_cos_raw = lambda x: jax.pure_callback(lambda a: np.cos(np.asarray(a)).astype(a.dtype),
jax.ShapeDtypeStruct(x.shape, x.dtype), x, vmap_method="expand_dims")
@jax.custom_jvp
def np_sin_v1(x):
return jax.pure_callback(lambda a: np.sin(np.asarray(a)).astype(a.dtype),
jax.ShapeDtypeStruct(x.shape, x.dtype), x, vmap_method="expand_dims")
np_sin_v1.defjvp(lambda p, t: (np_sin_v1(p[0]), np_cos_raw(p[0]) * t[0]))
xs = jnp.linspace(0.0, 1.0, 4)
np.testing.assert_allclose(jax.grad(lambda v: np_sin_v1(v).sum())(xs), np.cos(np.asarray(xs)), rtol=1e-6)
try:
jax.hessian(lambda v: np_sin_v1(v).sum())(xs)
ok = True
except ValueError as e:
ok = False
print("hessian(v1) → ValueError:", str(e)[:80])
assert not ok
# Fix: sin and cos are each other's derivatives, so give both a custom_jvp
@jax.custom_jvp
def np_sin(x):
return jax.pure_callback(lambda a: np.sin(np.asarray(a)).astype(a.dtype), jax.ShapeDtypeStruct(x.shape, x.dtype), x, vmap_method="expand_dims")
@jax.custom_jvp
def np_cos(x):
return np_cos_raw(x)
np_sin.defjvp(lambda p, t: (np_sin(p[0]), np_cos(p[0]) * t[0]))
np_cos.defjvp(lambda p, t: (np_cos(p[0]), -np_sin(p[0]) * t[0]))
H = jax.hessian(lambda v: np_sin(v).sum())(xs)
np.testing.assert_allclose(np.diag(H), -np.sin(np.asarray(xs)), rtol=1e-5)
print("Hessian diagonal after the fix:", np.diag(H))
# Check: a shape mismatch errors at run time; the callback doesn't run when its output is unused
try:
jax.jit(lambda v: jax.pure_callback(lambda a: np.asarray(a)[:2], jax.ShapeDtypeStruct(v.shape, v.dtype), v))(jnp.arange(3.0)).block_until_ready()
shape_err = None
except Exception as e:
shape_err = type(e).__name__
print("Wrong return shape →", shape_err)
assert shape_err == "JaxRuntimeError"
counter = []
def impure(a):
counter.append(1) # ❌ a side effect inside pure_callback (counterexample)
return np.asarray(a)
out = jax.jit(lambda v: (jax.pure_callback(impure, jax.ShapeDtypeStruct(v.shape, v.dtype), v), v + 1)[1])(jnp.arange(3.0))
out.block_until_ready()
print("Callback runs when the output is unused:", len(counter))
assert len(counter) == 0
# Timing: pure_callback (NumPy sin) vs. native jnp.sin (printed only)
xbig = jax.random.normal(jax.random.key(1), (1000, 1000))
f_cb = jax.jit(lambda v: jax.pure_callback(lambda a: np.sin(np.asarray(a)), jax.ShapeDtypeStruct(v.shape, v.dtype), v) * 2 + 1)
f_native = jax.jit(lambda v: jnp.sin(v) * 2 + 1)
np.testing.assert_allclose(f_cb(xbig), f_native(xbig), rtol=1e-5, atol=1e-6)
t_cb, t_nat = bench(f_cb, xbig), bench(f_native, xbig)
print(f"pure_callback: {t_cb:.3f} ms | native jnp: {t_nat:.3f} ms | ratio {t_cb / t_nat:.1f}x")
5.5 Introduction to JAX primitives¶
5.5 Introduction to JAX primitives¶
What a primitive is: a primitive is one of the most basic atomic operations JAX knows about. NumPy-style functions like jnp.add and jnp.dot are just Python wrappers; during tracing they are broken down into primitives and recorded in the jaxpr. Common primitives include add, mul, dot_general, reduce_sum, transpose, conv_general_dilated, and the control-flow primitives scan and cond.
Why they matter (the course gives three reasons):
- Transformations act on primitives:
jittraces a function into a jaxpr and hands it to XLA to compile;gradrelies on each primitive's differentiation rule;vmaprelies on each primitive's batching rule. Transformations never look at Python code, only at primitives. - The compilation interface: primitives connect the high-level API to the compiler. A jaxpr is lowered to compiler IR and then compiled to CPU/GPU/TPU machine code.
- Extensibility: to plug in a specialized algorithm written in C++/CUDA, or define an operation with an unusual differentiation rule, you define a new primitive and give it a full set of rules (5.6–5.9).
Primitives vs. functions: jnp.add is a Python function; during tracing it registers the primitive add in the jaxpr. All the rules that transformations and compilation need hang off the primitive object.
⚠️ Note: the control-flow primitives are named
scanandcond, notscan_pandcond_p; the primitive registered byjnp.addis also calledadd, notadd_p.*_pis just the Python variable name that holds the primitive object in the source code; the primitive's own name (.name, the one shown in the jaxpr) has no_p. Also, JAX's compilation pipeline today is jaxpr → StableHLO (an MLIR dialect) → XLA compilation, not jaxpr straight to XLA HLO; see the.lower().as_text()output below.
# See which primitives a function breaks down into
def f(x, y):
return jnp.sum(jnp.sin(x) * y + jnp.dot(x, y))
closed = jax.make_jaxpr(f)(jnp.ones(3), jnp.ones(3))
print(closed)
names = [eqn.primitive.name for eqn in closed.jaxpr.eqns]
print("Primitive sequence:", names)
assert {"sin", "mul", "dot_general", "add", "reduce_sum"} <= set(names)
# Names of the control-flow primitives
def g(x):
y = jax.lax.cond(x.sum() > 0, jnp.sin, jnp.cos, x)
return jax.lax.scan(lambda c, v: (c + v, c), 0.0, y)[0]
cf_names = {e.primitive.name for e in jax.make_jaxpr(g)(jnp.ones(3)).jaxpr.eqns}
print("Control-flow primitives:", cf_names & {"cond", "scan"})
assert {"cond", "scan"} <= cf_names
import jax.extend.core.primitives as P # semi-public entry point for built-in primitive objects
print("P.scan_p.name =", P.scan_p.name, "| P.cond_p.name =", P.cond_p.name, "| jax.lax.add_p.name =", jax.lax.add_p.name)
assert (P.scan_p.name, P.cond_p.name, jax.lax.add_p.name) == ("scan", "cond", "add")
# Primitives carry various rule tables: JVP, transpose, batching, lowering
from jax.interpreters import ad, batching, mlir
print("sin has a JVP rule :", P.sin_p in ad.primitive_jvps)
print("mul has a transpose rule :", P.mul_p in ad.primitive_transposes)
# Note: in 0.12, batching.primitive_batchers is a "write-only proxy" (only [p] = rule works); you can't query it with `in`.
# It wraps the rule and stores it in batching.fancy_primitive_batchers
try:
P.sin_p in batching.primitive_batchers
except TypeError as e:
print("primitive_batchers doesn't support `in`:", e)
print("sin has a batching rule :", P.sin_p in batching.fancy_primitive_batchers)
assert P.sin_p in ad.primitive_jvps and P.mul_p in ad.primitive_transposes and P.sin_p in batching.fancy_primitive_batchers
# The compilation pipeline: jaxpr → StableHLO (MLIR text)
hlo_text = jax.jit(f).lower(jnp.ones(3), jnp.ones(3)).as_text()
print(hlo_text)
assert "stablehlo.sine" in hlo_text and "stablehlo.dot_general" in hlo_text
🔍 Extra¶
⚪ Optional: when you need to define a custom primitive from scratch
- The rules a primitive needs to work fully (the next sections implement them one by one):
| Rule | Where it's registered (0.12) | Transformation it serves | Error when missing |
|---|---|---|---|
| impl (eager implementation) | p.def_impl | Direct calls (not under jit) | Evaluation rule ... not implemented |
| Abstract evaluation | p.def_abstract_eval | All tracing: jit / eval_shape / grad / vmap | Abstract evaluation ... not implemented |
| lowering | mlir.register_lowering | jit (compilation) | MLIR translation rule ... not found |
| JVP | ad.primitive_jvps[p] | jvp / grad | Differentiation rule ... not implemented |
| Transpose | ad.primitive_transposes[p] | grad / vjp (only linear primitives need it) | Transpose rule ... not implemented |
| batching | batching.primitive_batchers[p] | vmap | Batching rule ... not implemented |
- In 0.12,
batching.primitive_batchersis a write-only proxy:batching.primitive_batchers[p] = rulewraps the rule with a new signature and stores it inbatching.fancy_primitive_batchers. It doesn't supportinand can't be read; query the latter instead. Also, when none of the arguments is batched, the wrapper doesn't call your rule but callsbinddirectly. Semi-internal APIs change across versions, and this is a textbook example. - All these registries are semi-internal APIs. Most "I want to hook up external code" needs are covered by
pure_callbackorjax.ffi.ffi_callpluscustom_vjp(5.10); you don't need a custom primitive.
5.6 Defining custom primitives¶
5.6 Defining custom primitives¶
Motivation: jax.numpy and jax.lax already cover the vast majority of operations. You only need a new primitive in a few cases: ① performance, when you need to call a precompiled, highly optimized kernel; ② unsupported operations, which existing primitives can't express or which need special hardware features; ③ custom hardware. The course uses an analogy: primitives are the "vocabulary" JAX uses to describe computation, and a custom primitive adds a new word to that vocabulary.
Defining a primitive means spelling out four things: identity (a unique name), abstract behavior (deriving the output shape and dtype from only the input shapes and dtypes), execution (how each backend implements it, usually by lowering to compiler IR), and differentiation (JVP / transpose rules).
The course's five steps:
- Create the primitive object; its name shows up in the jaxpr.
- Write a wrapper function for users that calls
p.bind(*args, **params).binddoesn't compute the result right away; it hands the work to the current trace: in eager mode it calls the impl, under jit it records the call in the jaxpr. Static parameters that affect the output structure (such assome_parameter) must be passed tobindas keyword arguments. - Register the abstract evaluation rule with
def_abstract_eval(5.7). - Register the backend lowering rule with
mlir.register_lowering(5.8). - Register the differentiation rules (5.9).
The course also points out: if the primitive only calls external code, it's often enough to put jax.custom_vjp / jax.custom_jvp directly on the wrapper function, which is easier to maintain than defining rules on the primitive. Custom primitives come at a cost: high complexity, internal APIs that change across versions, and hard debugging. So consider jax.lax, callbacks, or a different formulation first.
This section walks through the whole process with the course's my_custom_op. For its math, we use the $f(x) = x^2 + c$ mentioned in the JVP comment ($c$ = some_parameter):
$$f(x; c) = x^2 + c,\qquad \dot y = 2x\,\dot x$$
⚠️ Note: in JAX 0.12,
from jax.core import PrimitiveraisesImportError, becausejax.coreno longer exportsPrimitive. The recommended path isfrom jax.extend.core import Primitive.ShapedArrayis still atjax.core.ShapedArray; it's not injax.extend.core.
# ❌ Old style (outdated/wrong): from jax.core import Primitive — jax.core no longer exports Primitive in 0.12
try:
from jax.core import Primitive as _P
except ImportError as e:
print("ImportError:", short_path(e))
import jax.core, jax.extend.core
print("jax.core has Primitive:", hasattr(jax.core, "Primitive"), "| has ShapedArray:", hasattr(jax.core, "ShapedArray"))
print("jax.extend.core has Primitive:", hasattr(jax.extend.core, "Primitive"), "| has ShapedArray:", hasattr(jax.extend.core, "ShapedArray"))
# ✅ Recommended: steps 1–3 — the primitive object, the user function, impl and abstract evaluation
from jax.extend.core import Primitive # semi-internal API: path verified on 0.12
from jax.core import ShapedArray # abstract value type
from jax.interpreters import mlir, ad, batching
my_custom_op_p = Primitive("my_custom_op") # step 1: a unique name
def my_custom_op(x, *, some_parameter): # step 2: the user-facing function
return my_custom_op_p.bind(x, some_parameter=some_parameter)
@my_custom_op_p.def_impl # eager implementation: receives concrete arrays
def _my_custom_op_impl(x, *, some_parameter):
return x * x + some_parameter
@my_custom_op_p.def_abstract_eval # step 3: looks only at shape/dtype, computes no values
def _my_custom_op_abstract_eval(abstract_x, *, some_parameter):
if some_parameter < 0:
raise ValueError("some_parameter must be non-negative")
return ShapedArray(abstract_x.shape, abstract_x.dtype)
x = jnp.array([1.0, 2.0, 3.0])
print("eager:", my_custom_op(x, some_parameter=1.0))
print(jax.make_jaxpr(partial(my_custom_op, some_parameter=1.0))(x))
np.testing.assert_array_equal(my_custom_op(x, some_parameter=1.0), [2.0, 5.0, 10.0])
# No lowering yet → jit fails; no JVP → grad fails; no batching → vmap fails
for name, fn in [("jit", lambda: jax.jit(partial(my_custom_op, some_parameter=1.0))(x)),
("grad", lambda: jax.grad(lambda v: my_custom_op(v, some_parameter=1.0).sum())(x)),
("vmap", lambda: jax.vmap(partial(my_custom_op, some_parameter=1.0))(jnp.ones((2, 3))))]:
try:
fn(); print(name, "→ unexpectedly succeeded")
except NotImplementedError as e:
print(f"{name:5s}→ NotImplementedError: {str(e).splitlines()[0][:80]}")
# ✅ Recommended: step 4 — lowering (details in 5.8; here mlir.lower_fun lowers a jnp implementation to StableHLO)
mlir.register_lowering(
my_custom_op_p,
mlir.lower_fun(lambda x, *, some_parameter: x * x + some_parameter, multiple_results=False))
y_jit = jax.jit(partial(my_custom_op, some_parameter=1.0))(x)
print("jit:", y_jit)
np.testing.assert_array_equal(y_jit, [2.0, 5.0, 10.0])
# ❌ Old style (outdated/wrong): building a zero tangent with ad.Zero.from_value(primal_out) in a JVP rule — this method doesn't exist in 0.12
print("ad.Zero has from_value:", hasattr(ad.Zero, "from_value"))
try:
ad.Zero.from_value(jnp.ones(3))
except AttributeError as e:
print("AttributeError:", e)
# ✅ Recommended: step 5 — the JVP rule (ẏ = 2x·ẋ); build zero tangents with ad.Zero(tangent aval)
def _my_custom_op_jvp(primals, tangents, *, some_parameter):
x, = primals
x_dot, = tangents
primal_out = my_custom_op_p.bind(x, some_parameter=some_parameter)
if type(x_dot) is ad.Zero: # the input is treated as a constant
tangent_out = ad.Zero(jax.typeof(primal_out).to_tangent_aval())
else:
tangent_out = 2 * x * x_dot # expressed with existing primitives, linear in x_dot
return primal_out, tangent_out
ad.primitive_jvps[my_custom_op_p] = _my_custom_op_jvp
g = jax.grad(lambda v: my_custom_op(v, some_parameter=1.0).sum())(x)
gj = jax.jit(jax.grad(lambda v: my_custom_op(v, some_parameter=1.0).sum()))(x)
print("grad:", g, "| jit(grad):", gj)
np.testing.assert_allclose(g, 2 * np.asarray(x)); np.testing.assert_allclose(gj, 2 * np.asarray(x))
# Batching for an elementwise primitive: batching.defvectorized does it in one line (the batch dim is passed straight through)
batching.defvectorized(my_custom_op_p)
xb = jnp.arange(6.0).reshape(2, 3)
yb = jax.vmap(partial(my_custom_op, some_parameter=1.0))(xb)
np.testing.assert_array_equal(yb, np.asarray(xb) ** 2 + 1)
vg = jax.vmap(jax.grad(lambda v: my_custom_op(v, some_parameter=0.0).sum()))(xb)
np.testing.assert_array_equal(vg, 2 * np.asarray(xb))
print("vmap / vmap(grad) passed")
⚠️ Note:
ad.primitive_transposesregisters a transpose rule, not a VJP. JAX's reverse mode = JVP linearization + transposition. You only need a transpose rule when the tangent part of the JVP uses this primitive itself (in which case the primitive is linear in the tangent). The JVP above expresses the tangent as2 * x * x_dot(the built-inmul), andmulcomes with its own transpose rule, so registering only the JVP is enough forjax.gradto work, as the asserts above confirm. Section 5.9 covers the cases where you have to write a transpose rule yourself.
5.7 Implementing abstract evaluation rules¶
5.7 Implementing abstract evaluation rules¶
What abstract evaluation is: before jit, autodiff or vmap actually compute anything, JAX first needs to know the shape and dtype of each operation's output, without computing concrete values. It traces the program with abstract values (mainly jax.core.ShapedArray, which holds only shape and dtype).
Purpose: abstract evaluation is the primitive's "function signature at the level of types and shapes". When tracing hits a primitive, JAX calls it with (abstract inputs, params) and gets abstract outputs. The course gives three reasons:
- Static analysis: shape mismatches and type errors are caught before anything is compiled or run.
- Transformation compatibility: vmap relies on it to derive batched shapes, and grad relies on it to build the structure of the backward pass.
- Compilation: XLA needs the exact shape and type of every operation to generate code.
Parameters (params): if the output shape depends on non-array parameters (convolution strides, reduction axes and so on), these are passed in as keyword arguments to bind, and the abstract evaluation rule gets them from its keyword arguments too.
Where it fits: abstract evaluation is the first rule to write for a new primitive; both lowering and JVP use the shape information it provides. The jaxpr is held together by these abstract values.
The course gives two examples: custom_add_one (adds 1, and checks the number of inputs, their type and that the dtype is floating point), and custom_reduce_sum (removes dimensions according to an axis parameter). Below, both are run in full.
# ✅ Course example 1: abstract evaluation for custom_add_one (with input checks)
custom_add_one_p = Primitive("custom_add_one")
@custom_add_one_p.def_abstract_eval
def custom_add_one_abstract_eval(*avals, **params):
if len(avals) != 1:
raise ValueError("custom_add_one needs exactly one input operand.")
input_aval = avals[0]
if not isinstance(input_aval, ShapedArray):
raise TypeError(f"Expected a ShapedArray, got {type(input_aval)}")
if not jnp.issubdtype(input_aval.dtype, jnp.floating):
raise TypeError(f"custom_add_one only supports floating-point inputs, got {input_aval.dtype}")
return ShapedArray(input_aval.shape, input_aval.dtype)
impl_calls = []
@custom_add_one_p.def_impl
def _add_one_impl(x):
impl_calls.append(1)
return x + 1.0
# jax.eval_shape only does abstract evaluation: you get the shape, and the impl never runs
out_sds = jax.eval_shape(custom_add_one_p.bind, jax.ShapeDtypeStruct((4, 5), jnp.float32))
print("eval_shape:", out_sds, "| impl calls:", len(impl_calls))
assert out_sds.shape == (4, 5) and out_sds.dtype == jnp.float32 and impl_calls == []
# Type errors are caught at tracing time (before compiling or running anything)
try:
jax.eval_shape(custom_add_one_p.bind, jax.ShapeDtypeStruct((3,), jnp.int32))
caught = False
except TypeError as e:
caught = True
print("TypeError at tracing time:", e)
assert caught
# ✅ Course example 2: custom_reduce_sum — the output shape is decided by the axis parameter
custom_reduce_sum_p = Primitive("custom_reduce_sum")
def custom_reduce_sum(x, axis):
axis = tuple(sorted(a % x.ndim for a in (axis if isinstance(axis, tuple) else (axis,))))
return custom_reduce_sum_p.bind(x, axis=axis) # axis must be hashable (a tuple)
@custom_reduce_sum_p.def_abstract_eval
def custom_reduce_sum_abstract_eval(aval, *, axis):
if not isinstance(aval, ShapedArray):
raise TypeError("The input must be a ShapedArray")
output_shape = tuple(d for i, d in enumerate(aval.shape) if i not in axis)
return ShapedArray(output_shape, aval.dtype)
custom_reduce_sum_p.def_impl(lambda x, *, axis: jnp.sum(x, axis=axis))
mlir.register_lowering(custom_reduce_sum_p,
mlir.lower_fun(lambda x, *, axis: jnp.sum(x, axis=axis), multiple_results=False))
my_array = jnp.arange(24.0).reshape(2, 3, 4)
print(jax.eval_shape(lambda a: custom_reduce_sum(a, (0,)), my_array))
assert jax.eval_shape(lambda a: custom_reduce_sum(a, (0, 2)), my_array).shape == (3,)
r = jax.jit(lambda a: custom_reduce_sum(a, (0,)))(my_array)
np.testing.assert_array_equal(r, np.asarray(my_array).sum(0))
print(jax.make_jaxpr(lambda a: custom_reduce_sum(a, -1))(my_array))
⚠️ Note: for a primitive with multiple outputs, having abstract evaluation return a tuple of
ShapedArrays isn't enough; you must also setp.multiple_results = Truefirst. Otherwise, returning a tuple from abstract evaluation raisesValueError: ... should return a tuple or a list iff ... multiple_results. The impl and lowering must also be written for multiple outputs (lower_fun(..., multiple_results=True)).
# A multi-output primitive: minmax(x) -> (min, max)
minmax_bad = Primitive("minmax_bad") # forgot to set multiple_results
minmax_bad.def_abstract_eval(lambda a: (ShapedArray((), a.dtype), ShapedArray((), a.dtype)))
try:
jax.eval_shape(minmax_bad.bind, jnp.arange(5.0))
err = None
except ValueError as e:
err = e; print("ValueError:", str(e)[:110])
assert err is not None
minmax_p = Primitive("minmax")
minmax_p.multiple_results = True
minmax_p.def_impl(lambda x: (jnp.min(x), jnp.max(x)))
minmax_p.def_abstract_eval(lambda a: (ShapedArray((), a.dtype), ShapedArray((), a.dtype)))
mlir.register_lowering(minmax_p, mlir.lower_fun(lambda x: (jnp.min(x), jnp.max(x)), multiple_results=True))
lo, hi = jax.jit(minmax_p.bind)(jnp.array([3.0, -1.0, 7.0]))
print("min, max =", lo, hi)
assert (float(lo), float(hi)) == (-1.0, 7.0)
🔍 Extra¶
⚪ Optional: when you write an abstract evaluation rule for a custom primitive
- Abstract evaluation can only contain shape logic: it receives
ShapedArrays, not arrays. Anything that reads a value inside it, likefloat(x)ornp.asarray(x), fails. Anything that depends on a "value" (such as the output shape) can only come from static parameters. - Parameters must be hashable:
bind's keyword arguments become part of the jaxpr equation and also take part in jit caching, so use tuples, not lists or arrays. Thecustom_reduce_sumabove first normalizesaxisinto a sorted tuple, so equivalent ways of writing it hit the same cache entry. jax.eval_shapeis great for debugging abstract evaluation: it only traces and never executes, so you can check shape inference and error messages quickly.
5.8 Implementing lowering rules for backends (CPU/GPU/TPU)¶
5.8 Implementing lowering rules for backends¶
Why you need it: jit hands the function to XLA to compile. When it meets a primitive, JAX has to translate it into low-level instructions the compiler understands. Without this "lowering rule", XLA doesn't know what code to generate for the custom operation. The best implementation may differ between backends, so you can register rules per platform.
The mechanism described in the course: a translation rule receives XLA values, builds an HLO graph with the XLA builder (builder.ConstantF32Scalar, builder.Add, builder.Mul, builder.DotGeneral, builder.Reshape…), and is registered with xla_client.register_translation(prim, rule, platform=...).
Backend considerations: on CPU, XLA can use multiple cores and vector instructions (AVX and so on). On GPU, if there's no matching HLO op, you can embed a CUDA kernel through CustomCall. On TPU you need to think about the matrix unit (MXU), and usually organize the computation as large matrix multiplications.
The course example: ReLU $\mathrm{relu}(x) = \max(0, x)$. First get the shape and dtype from the operand, build a zero tensor of the same shape, then take an elementwise Max. XLA knows how to compile Max efficiently, so one rule is enough for all backends.
More complex cases: combining several existing HLO ops; writing a dedicated GPU kernel and calling it through CustomCall; understanding how XLA fusion affects the generated code. Debugging usually means looking at the generated IR.
⚠️ Note:
from jax.lib import xla_client,xla_client.register_translationand the XLA builder (builder.Add,builder.GetShapeand so on) all belong to the old XLA translation-rule API. In JAX 0.12, importingjax.lib.xla_clientfails, andjax.interpreters.xlahas noregister_translation(only a few members likeapply_primitiveremain).The current way is MLIR lowering: register with
jax.interpreters.mlir.register_lowering(prim, rule, platform=None|'cpu'|'cuda'|'rocm'|'tpu'). The rule's signature isrule(ctx, *mlir_values, **params) -> list[ir.Value], and it produces StableHLO. The simplest way ismlir.lower_fun(<jnp implementation>, multiple_results=...). For fine control, write ops by hand withjax.extend.mlir.dialects.stablehlo. To call external C++, usejax.ffi.ffi_lowering(5.10, Exercise 4).
# ❌ Old style (outdated/wrong): xla_client.register_translation + XLA builder — gone in 0.12
try:
from jax.lib import xla_client
except ImportError as e:
print("ImportError:", short_path(e))
from jax.interpreters import xla
print("Public members of jax.interpreters.xla:", [a for a in dir(xla) if not a.startswith("_")])
print("Has register_translation:", hasattr(xla, "register_translation"))
assert not hasattr(xla, "register_translation")
# ✅ Recommended A: mlir.lower_fun — write the semantics in jnp and let JAX lower it to StableHLO
my_relu_p = Primitive("my_relu")
my_relu_p.def_impl(lambda x: jnp.maximum(x, 0))
my_relu_p.def_abstract_eval(lambda a: ShapedArray(a.shape, a.dtype))
# Generic rule (all platforms): implemented with where
mlir.register_lowering(my_relu_p, mlir.lower_fun(lambda x: jnp.where(x > 0, x, jnp.zeros_like(x)), multiple_results=False))
txt_generic = jax.jit(my_relu_p.bind).lower(jnp.ones(3)).as_text()
print(txt_generic)
assert "stablehlo.select" in txt_generic
# ✅ Recommended B: hand-written StableHLO (the course's ReLU example), registered only for the 'cpu' platform
from jax.extend.mlir.dialects import stablehlo as hlo
def my_relu_lowering_cpu(ctx, x_operand):
aval_out, = ctx.avals_out # output aval from abstract evaluation
out_type = mlir.aval_to_ir_type(ctx.module_context, aval_out) # 0.12: needs module_context
zero = mlir.ir_constant(np.zeros((), aval_out.dtype)) # scalar 0
zeros = hlo.broadcast_in_dim(out_type, zero, mlir.dense_int_array([])) # broadcast to the same shape
return [hlo.maximum(zeros, x_operand)] # must return a list of Values
mlir.register_lowering(my_relu_p, my_relu_lowering_cpu, platform="cpu") # platform-specific rules take priority over the generic one
xr = jnp.array([-1.0, 2.0, -3.0, 4.0])
txt_cpu = jax.jit(my_relu_p.bind).lower(xr).as_text()
print(txt_cpu)
assert "stablehlo.maximum" in txt_cpu and "stablehlo.select" not in txt_cpu
np.testing.assert_array_equal(jax.jit(my_relu_p.bind)(xr), [0.0, 2.0, 0.0, 4.0])
🔍 Extra¶
⚪ Optional: when you hand-write a lowering, or register different implementations per platform
- Values of
platform:Nonemeans a generic rule; platform-specific rules ('cpu','cuda','rocm','tpu') take priority. On GPU the platform names are'cuda'/'rocm', not'gpu'. - Lowering is not impl: jit uses the lowering, eager mode uses
def_impl. If the two implementations disagree, the same function gives different results with and without jit, and that kind of bug is hard to spot. An easy fix is to have the impl reuse the lowering:p.def_impl(partial(xla.apply_primitive, p)), which jits just this one primitive on the fly for eager calls (Exercise 4 uses this). - Useful fields on
ctx:ctx.avals_in/ctx.avals_out(abstract values),ctx.module_context(needed to build types),ctx.platforms.
5.9 Differentiation rules for custom primitives¶
5.9 Differentiation rules for custom primitives¶
Why you need them: a primitive with only abstract evaluation and lowering can't take part in autodiff yet. Calling jax.grad, jax.jvp or jax.vjp on it raises an error, because JAX doesn't know how to propagate derivatives through it.
Forward mode (JVP): let $y = f(x)$ with Jacobian $J = \partial f / \partial x$. The JVP computes
$$\dot y = J\,\dot x .$$
The rule takes (primals, tangents) and returns (primal_out, tangent_out). A tangent has the same shape and dtype as its primal. For an input treated as a constant, the tangent is the symbolic zero ad.Zero.
Reverse mode (VJP): computes
$$\bar x^{\top} = \bar y^{\top} J .$$
jax.grad uses reverse mode. The VJP described in the course has the structure "save residuals in the forward pass, compute input cotangents in backward_pass", which is how custom_vjp is written.
The course example custom_scale: $y = x \cdot c$ ($c$ is a scalar factor).
$$\dot y = \dot x\, c + x\, \dot c,\qquad \bar x = \bar y\, c,\qquad \bar c = \sum_i \bar y_i\, x_i$$
$c$ is a scalar, so $\bar c$ sums over all elements. We check this with $L = \sum (x c)^2$:
$$\frac{\partial L}{\partial x} = 2 x c^2 = [8, 16, 24],\qquad \frac{\partial L}{\partial c} = 2c\sum x^2 = 56\qquad (x=[1,2,3],\ c=2)$$
Expected JVP value: with $\dot x = [0.1, 0.2, 0.3]$ and $\dot c = 0.5$, $\dot y = [0.7, 1.4, 2.1]$.
How to verify (from the course): ① a numerical check, using jax.test_util.check_grads to compare finite differences against the custom rules; ② check that JVP and VJP agree.
⚠️ Note: older material often registers rules with
custom_scale_p.def_jvp/def_vjp, callsjax.core.get_avalinside abstract evaluation, and builds zeros withad.Zero.zero_if_zero. In JAX 0.12,Primitivehas nodef_jvp/def_vjpmethods at all, and neitherjax.core.get_avalnorad.Zero.zero_if_zeroexists. The correct approach at the primitive level is:
- JVP:
ad.primitive_jvps[p] = jvp_rule;- Reverse mode: don't register a VJP. JAX linearizes with the JVP first, then transposes the tangent computation. If the tangent part of the JVP uses this primitive, register
ad.primitive_transposes[p] = transpose_rule;- The "forward + residuals + backward_pass" structure belongs to
jax.custom_vjp, which fits at the wrapper-function level (see 5.10).One more thing: if you write
custom_scale_p.bind(x, factor=factor), factor becomes a parameter, while the abstract evaluation signature(x_abs, factor_abs)treats it as a second operand, so the two don't match. The right way: if factor needs a gradient, it must be passed tobindas an operand (a positional argument). Parameters (params) are static and can't be differentiated.
# ❌ Old style (outdated/wrong): Primitive.def_jvp / def_vjp, jax.core.get_aval, ad.Zero.zero_if_zero don't exist; bind's params contradict the abstract eval signature
print("Primitive has def_jvp:", hasattr(Primitive, "def_jvp"), "| has def_vjp:", hasattr(Primitive, "def_vjp"))
print("jax.core has get_aval:", hasattr(jax.core, "get_aval"), "| ad.Zero has zero_if_zero:", hasattr(ad.Zero, "zero_if_zero"))
print("jax.ad has Zero:", hasattr(jax.ad, "Zero"), "(jax.ad is now a different, new namespace)")
course_scale_p = Primitive("course_scale")
course_scale_p.def_impl(lambda x, factor: x * factor)
course_scale_p.def_abstract_eval(lambda x_abs, factor_abs: ShapedArray(x_abs.shape, x_abs.dtype))
try:
course_scale_p.def_jvp(lambda p, t: None)
except AttributeError as e:
print("def_jvp → AttributeError:", e)
try:
jax.jit(lambda x: course_scale_p.bind(x, factor=2.0))(jnp.ones(3)) # factor passed as a parameter
except TypeError as e:
print("bind(x, factor=...) → TypeError:", str(e).splitlines()[0][:120])
# ✅ Recommended: the custom_scale primitive — both x and factor are operands
custom_scale_p = Primitive("custom_scale")
def custom_scale(x, factor):
return custom_scale_p.bind(x, factor)
@custom_scale_p.def_impl
def _custom_scale_impl(x, factor):
return x * factor
@custom_scale_p.def_abstract_eval
def _custom_scale_abstract_eval(x_aval, factor_aval):
if factor_aval.shape != ():
raise TypeError(f"factor must be a scalar, got shape {factor_aval.shape}")
return ShapedArray(x_aval.shape, x_aval.dtype)
mlir.register_lowering(custom_scale_p, mlir.lower_fun(lambda x, f: x * f, multiple_results=False))
# JVP: ẏ = custom_scale(ẋ, c) + x·ċ — the first term uses this primitive itself (linear in ẋ), so reverse mode needs a transpose rule
def _custom_scale_jvp(primals, tangents):
x, factor = primals
x_dot, factor_dot = tangents
y = custom_scale(x, factor)
terms = []
if type(x_dot) is not ad.Zero:
terms.append(custom_scale(x_dot, factor))
if type(factor_dot) is not ad.Zero:
terms.append(x * factor_dot)
if not terms:
return y, ad.Zero(jax.typeof(y).to_tangent_aval())
return y, sum(terms[1:], terms[0])
ad.primitive_jvps[custom_scale_p] = _custom_scale_jvp
x_val = jnp.array([1.0, 2.0, 3.0]); factor_val = jnp.float32(2.0)
primal_out, tangent_out = jax.jvp(custom_scale, (x_val, factor_val), (jnp.array([0.1, 0.2, 0.3]), jnp.float32(0.5)))
print("JVP:", primal_out, tangent_out)
np.testing.assert_allclose(tangent_out, [0.7, 1.4, 2.1], rtol=1e-6)
def loss_fn(x, factor):
y = custom_scale(x, factor)
return jnp.sum(y * y)
# Gradient w.r.t. factor: the tangent path only goes through the built-in mul, so no transpose rule is needed → already works
print("dL/dfactor (no transpose rule yet):", jax.grad(loss_fn, argnums=1)(x_val, factor_val))
assert float(jax.grad(loss_fn, argnums=1)(x_val, factor_val)) == 56.0
# Gradient w.r.t. x: needs custom_scale(ẋ, c) to be transposed
try:
jax.grad(loss_fn, argnums=0)(x_val, factor_val)
except NotImplementedError as e:
print("dL/dx →", str(e).splitlines()[0][:90])
# ✅ Transpose rule: for the linear argument x (an UndefinedPrimal) return the cotangent x̄ = ȳ·c; return None for the nonlinear argument factor
def _custom_scale_transpose(ct, x, factor):
assert ad.is_undefined_primal(x) and not ad.is_undefined_primal(factor)
if type(ct) is ad.Zero:
ct_x = ad.Zero(x.aval)
else:
ct_x = custom_scale(ct, factor) # ȳ·c is still expressed with this primitive
return ct_x, None # one entry per operand
ad.primitive_transposes[custom_scale_p] = _custom_scale_transpose
grad_x = jax.grad(loss_fn, argnums=0)(x_val, factor_val)
grad_factor = jax.grad(loss_fn, argnums=1)(x_val, factor_val)
print("dL/dx =", grad_x, "| dL/dfactor =", grad_factor)
np.testing.assert_allclose(grad_x, [8.0, 16.0, 24.0]); assert float(grad_factor) == 56.0
gx, gf = jax.jit(jax.grad(loss_fn, argnums=(0, 1)))(x_val, factor_val)
np.testing.assert_allclose(gx, [8.0, 16.0, 24.0]); assert float(gf) == 56.0
# The course's verification method: check_grads (first and second order, forward and reverse)
from jax.test_util import check_grads
check_grads(custom_scale, (x_val, factor_val), order=2, modes=["fwd", "rev"], eps=1e-2)
print("check_grads passed.")
🔍 Extra: batching rules¶
⚪ Optional: when a custom primitive needs to support
vmap
To make a primitive work with vmap, you also register batching.primitive_batchers[p] = rule. The rule's signature is rule(batched_args, batch_dims, **params) -> (out, out_dim). Here batch_dims[i] is the position of the batch dimension of argument i, and None means that argument isn't batched. The rule's job is to process the whole batch in one go with a "computation without vmap", then tell JAX where the output's batch dimension is.
custom_scale has two cases:
- factor is not batched: $y = x\cdot c$ is elementwise anyway, so just call
custom_scale(x, c); the output batch dimension is the same as x's. - factor is batched (each example has its own $c_b$): move the batch dimension to the front and compute $y_b = x_b \cdot c_b$ with a broadcast multiply; the output batch dimension is 0.
def _custom_scale_batch(batched_args, batch_dims):
x, factor = batched_args
xd, fd = batch_dims
if fd is None: # factor isn't batched
return custom_scale(x, factor), xd
B = factor.shape[fd]
factor = jnp.moveaxis(factor, fd, 0) # (B,)
x = jnp.moveaxis(x, xd, 0) if xd is not None else jnp.broadcast_to(x, (B,) + x.shape)
return x * factor.reshape((B,) + (1,) * (x.ndim - 1)), 0
batching.primitive_batchers[custom_scale_p] = _custom_scale_batch
xb = jnp.arange(12.0).reshape(4, 3)
cb = jnp.arange(4.0)
np.testing.assert_array_equal(jax.vmap(custom_scale, (0, None))(xb, factor_val), np.asarray(xb) * 2)
np.testing.assert_array_equal(jax.vmap(custom_scale, (0, 0))(xb, cb), np.asarray(xb) * np.asarray(cb)[:, None])
np.testing.assert_array_equal(jax.vmap(custom_scale, (1, 0))(xb.T, cb), np.asarray(xb) * np.asarray(cb)[:, None])
np.testing.assert_array_equal(jax.vmap(custom_scale, (None, 0))(x_val, cb), np.outer(np.asarray(cb), np.asarray(x_val)))
# Composing vmap ∘ grad ∘ jit: per example, dL/dx = 2·x·c²
vg = jax.jit(jax.vmap(jax.grad(loss_fn)))(xb, cb)
np.testing.assert_allclose(vg, 2 * np.asarray(xb) * np.asarray(cb)[:, None] ** 2)
print("custom_scale passes jit / jvp / grad / vmap / vmap(grad) / check_grads")
print(jax.make_jaxpr(jax.grad(loss_fn))(x_val, factor_val))
🔍 Extra¶
⚪ Optional: when you write differentiation rules for a custom primitive
- Why JVP + transpose is all you need: JAX's reverse mode is
vjp = transpose(linearize(jvp)). As long as each primitive has a "forward derivative" and a "transpose of its linear part", forward mode, reverse mode and higher-order derivatives (such ascheck_grads(order=2)) can all be assembled automatically. This is a design choice at the primitive level, and it composes more easily than writing a separate VJP for every primitive. - Arguments of a transpose rule: linear inputs show up as
ad.UndefinedPrimal; test for them withad.is_undefined_primal, and their.avalgives the shape. Nonlinear inputs are ordinary values. The return value has one entry per operand, withNonein the positions of nonlinear inputs. - Zero tangents:
ad.Zero(aval)is a symbolic zero and takes no memory. In a rule, checktype(t) is ad.Zerofirst, and don't do arithmetic on it directly. When you need a real array of zeros, usead.instantiate_zeros(t). - When not to use a primitive: if all you want is a gradient for external code (a callback or FFI), wrapping the user function with
custom_jvp/custom_vjpis enough (5.4, 5.10), and you never need to touchad.primitive_*.
5.10 Practice: integrating a C++ function¶
5.10 Practice: integrating a C++ function¶
JAX/XLA is already highly optimized, but sometimes you still need to reuse an existing C++ library, or implement a performance-critical low-level operation yourself. The course presents two approaches: simple callback integration and a full custom primitive. The hands-on example uses the former (ctypes + pure_callback); the latter is only outlined.
Scenario: an elementwise function (pretend it's a complicated piece of legacy C++)
$$f(x) = x^2 + 10,\qquad f'(x) = 2x$$
The course's steps:
- C++ implementation: use
extern "C"to prevent C++ name mangling, so ctypes can find the symbol by name; computeinput[i]*input[i] + 10.0elementwise over adoublearray. - Compile to a shared library:
g++ -shared -fPIC -o custom_op.so custom_op.cppon Linux,g++ -shared -o custom_op.dylib custom_op.cppon macOS, and the matching MinGW or MSVC command on Windows. - ctypes wrapper: load the library with
ctypes.CDLLand setargtypes/restype. A NumPy wrapper function converts to float64, makes sure the memory is contiguous, allocates the output, and passes pointers to the C function. The expected result is[1,2,3] → [11,14,19]. pure_callbackintegration: this C++ function is mathematically pure, so it's a good fit for pure_callback. The course tests it eagerly and under jit (expecting[1,2,3,4] → [11,14,19,26]), and points out that jax.grad fails, because autodiff can't see into C++ code.- The alternative (outline only): define a Primitive → abstract evaluation → per-backend lowering (mentioning
xla_client.ops.CustomCallandxla.register_translation) → JVP/transpose → user function callingbind. - Summary: callbacks suit pure functions where you don't need gradients through the external code. Custom primitives give the tightest integration and support autodiff, but are complex.
This section first gets the ctypes approach working in 0.12 style, following the course's outline. Then it uses the currently recommended XLA FFI (jax.ffi) to plug in a genuinely compiled C++ RMSNorm kernel, and adds gradient and vmap support.
⚠️ Note: if you write this integration code the old way, 4 things break on 0.12:
from jax.experimental import host_callback as hcbraisesImportErrorright away. Even if it's never used after the import, the whole script fails at import time;from jax.experimental import pure_callbackraisesImportError; the correct entry point isjax.pure_callback;pure_callback(..., vectorized=True):vectorizedis no longer a pure_callback parameter. It gets forwarded to the callback as a kwarg and fails at run time. Usevmap_method="expand_dims"or"broadcast_all"instead;- With x64 off (the default),
jnp.arange(1.0, 5.0, dtype=jnp.float64)is truncated to float32 (with a UserWarning), so you can't treat it as float64.
xla_client.ops.CustomCallandxla.register_translation(and the earlier popular approach of exporting aPyCapsulewith pybind11) are the old way of hooking in custom calls. 0.12 recommends XLA FFI: on the C++ side, define a type-safe handler withXLA_FFI_DEFINE_HANDLER_SYMBOLfromxla/ffi/api/ffi.h; on the Python side, register it withjax.ffi.pycapsule+jax.ffi.register_ffi_target, then call it withjax.ffi.ffi_call. No pybind11 needed.
%%writefile cpp/custom_op.cpp
// Step 1: elementwise f(x) = x^2 + 10; extern "C" prevents name mangling so ctypes can call it
#include <cstdint>
extern "C" {
void custom_elementwise_func(const double* input, double* output, int size) {
for (int i = 0; i < size; ++i) {
output[i] = input[i] * input[i] + 10.0;
}
}
}
# Course step 2: compile to a shared library (local clang++; run via subprocess in the cell, and print errors as-is on failure)
def compile_so(src, out, extra=()):
cmd = [CXX, "-std=c++17", "-O2", "-shared", "-fPIC", *extra, "-o", str(out), str(src)]
print("$", short_path(" ".join(cmd)))
r = subprocess.run(cmd, capture_output=True, text=True)
if r.returncode != 0:
print(r.stdout, r.stderr)
raise RuntimeError(f"Compilation failed: {src}")
print("Compiled →", out.name, f"({out.stat().st_size} bytes)")
return out
CXX = shutil.which("clang++") or shutil.which("g++") or "c++" # clang++ on macOS; usually g++ on Linux / Colab
print(subprocess.run([CXX, "--version"], capture_output=True, text=True).stdout.splitlines()[0])
so_custom = compile_so(CPP_DIR / "custom_op.cpp", CPP_DIR / "libcustom_op.so")
assert so_custom.exists()
# Course step 3: ctypes wrapper (absolute path for the library)
lib = ctypes.CDLL(str(so_custom))
lib.custom_elementwise_func.argtypes = [ctypes.POINTER(ctypes.c_double),
ctypes.POINTER(ctypes.c_double),
ctypes.c_int]
lib.custom_elementwise_func.restype = None
def custom_op_numpy(x_np):
x_np = np.ascontiguousarray(np.asarray(x_np), dtype=np.float64) # float64 + contiguous memory
output_np = np.empty_like(x_np)
lib.custom_elementwise_func(x_np.ctypes.data_as(ctypes.POINTER(ctypes.c_double)),
output_np.ctypes.data_as(ctypes.POINTER(ctypes.c_double)),
x_np.size)
return output_np
result_np = custom_op_numpy(np.array([1.0, 2.0, 3.0]))
print("NumPy wrapper test:", result_np)
np.testing.assert_array_equal(result_np, [11.0, 14.0, 19.0])
np.testing.assert_array_equal(custom_op_numpy(np.arange(6.0).reshape(2, 3)), np.arange(6.0).reshape(2, 3) ** 2 + 10)
# ❌ Old style (outdated/wrong): host_callback / jax.experimental.pure_callback imports fail; float64 is truncated; vectorized=True is forwarded to the callback
for stmt in ["from jax.experimental import host_callback as hcb",
"from jax.experimental import pure_callback"]:
try:
exec(stmt); print(stmt, "→ succeeded")
except ImportError as e:
print(stmt, "→ ImportError")
with warnings.catch_warnings(record=True) as w:
warnings.simplefilter("always")
x_jax = jnp.arange(1.0, 5.0, dtype=jnp.float64)
print("Requested float64, actual dtype =", x_jax.dtype, "| warning:", str(w[0].message)[:70] if w else None)
import logging
_cb_logger = logging.getLogger("jax._src.callback") # when a callback fails JAX logs the full traceback; turn that off here
_cb_logger.disabled = True
try:
y = jax.pure_callback(custom_op_numpy, jax.ShapeDtypeStruct(x_jax.shape, x_jax.dtype), x_jax, vectorized=True)
y.block_until_ready()
except Exception as e:
print("vectorized=True →", type(e).__name__, "| root cause: custom_op_numpy() got an unexpected keyword argument 'vectorized'")
_cb_logger.disabled = False
# ✅ Recommended: jax.pure_callback + vmap_method + custom_jvp (f'(x) = 2x)
@jax.custom_jvp
def custom_op_jax(x):
return jax.pure_callback(
lambda a: custom_op_numpy(a).astype(a.dtype), # back to the declared dtype
jax.ShapeDtypeStruct(x.shape, x.dtype), x,
vmap_method="expand_dims") # elementwise op: one call for the whole batch is enough
@custom_op_jax.defjvp
def _custom_op_jvp(primals, tangents):
(x,), (x_dot,) = primals, tangents
return custom_op_jax(x), 2.0 * x * x_dot
x_jax = jnp.arange(1.0, 5.0) # float32
print("eager:", custom_op_jax(x_jax), "| jit:", jax.jit(custom_op_jax)(x_jax))
np.testing.assert_array_equal(jax.jit(custom_op_jax)(x_jax), [11.0, 14.0, 19.0, 26.0])
g = jax.grad(lambda v: jnp.sum(custom_op_jax(v)))(x_jax)
print("grad:", g)
np.testing.assert_array_equal(g, 2 * np.arange(1.0, 5.0))
np.testing.assert_array_equal(jax.vmap(custom_op_jax)(x_jax.reshape(2, 2)), (np.arange(1.0, 5.0) ** 2 + 10).reshape(2, 2))
The modern way: plugging in a C++ RMSNorm with XLA FFI (jax.ffi)¶
RMSNorm normalizes over the last dimension ($N$ is the length of the last dimension):
$$s = \Big(\tfrac{1}{N}\sum_{i} x_i^2 + \varepsilon\Big)^{-1/2},\qquad y_i = s\,x_i$$
For backpropagation ($\bar y$ is the output cotangent), $\partial s/\partial x_j = -s^3 x_j / N$ gives:
$$\bar x_j = s\,\bar y_j - \frac{s^3}{N}\, x_j \sum_i x_i\,\bar y_i$$
The FFI workflow:
- C++ side: write an ordinary function
ffi::Error Impl(attributes..., ffi::Buffer<T> inputs..., ffi::ResultBuffer<T> outputs...), then declare its signature withXLA_FFI_DEFINE_HANDLER_SYMBOL(symbol_name, Impl, ffi::Ffi::Bind().Attr<float>("eps").Arg<...>().Ret<...>()). At call time XLA checks types and passes shape information (dimensions()) along to C++. - Compile: the headers live under
jax.ffi.include_dir(). You only need the headers; no library to link. - Python side: load with
ctypes.cdll.LoadLibrary→ wrap the function pointer withjax.ffi.pycapsule(lib.<symbol>)→ register withjax.ffi.register_ffi_target(<name>, <capsule>, platform="cpu"). - Call:
jax.ffi.ffi_call(<name>, <output ShapeDtypeStruct>, vmap_method=...)(*inputs, **attributes). It produces a StableHLOcustom_call, so it works with jit out of the box, and there's no round trip through Python like with callbacks. - Gradients:
ffi_callisn't differentiable, so usejax.custom_vjpwith a forward kernel (which also outputs a residual $s$) and a backward kernel.
%%writefile cpp/rms_norm.cc
// RMSNorm with XLA FFI: three handlers — forward, forward with residuals, and backward
#include <cmath>
#include <cstdint>
#include <utility>
#include "xla/ffi/api/ffi.h"
namespace ffi = xla::ffi;
// RMSNorm over one row of length n; returns the scale factor s
static float ComputeRmsNorm(float eps, int64_t n, const float* x, float* y) {
float sm = 0.0f;
for (int64_t i = 0; i < n; ++i) sm += x[i] * x[i];
float scale = 1.0f / std::sqrt(sm / float(n) + eps);
for (int64_t i = 0; i < n; ++i) y[i] = x[i] * scale;
return scale;
}
// Backward: ct_x = s*ct_y - s^3/n * x * sum(x*ct_y)
static void ComputeRmsNormBwd(int64_t n, float scale, const float* x, const float* ct_y, float* ct_x) {
float dot = 0.0f;
for (int64_t i = 0; i < n; ++i) dot += x[i] * ct_y[i];
float f = scale * scale * scale * dot / float(n);
for (int64_t i = 0; i < n; ++i) ct_x[i] = scale * ct_y[i] - f * x[i];
}
// Returns (total element count, length of the last dim); all leading dims are treated as "rows"
template <ffi::DataType T>
static std::pair<int64_t, int64_t> GetDims(const ffi::Buffer<T>& b) {
auto dims = b.dimensions();
if (dims.size() == 0) return {0, 0};
return {b.element_count(), dims.back()};
}
static ffi::Error RmsNormImpl(float eps, ffi::Buffer<ffi::F32> x, ffi::ResultBuffer<ffi::F32> y) {
auto [total, last] = GetDims(x);
if (last == 0) return ffi::Error::InvalidArgument("RmsNorm input must be an array");
for (int64_t n = 0; n < total; n += last)
ComputeRmsNorm(eps, last, &(x.typed_data()[n]), &(y->typed_data()[n]));
return ffi::Error::Success();
}
static ffi::Error RmsNormFwdImpl(float eps, ffi::Buffer<ffi::F32> x, ffi::ResultBuffer<ffi::F32> y,
ffi::ResultBuffer<ffi::F32> res) {
auto [total, last] = GetDims(x);
if (last == 0) return ffi::Error::InvalidArgument("RmsNormFwd input must be an array");
for (int64_t n = 0, row = 0; n < total; n += last, ++row)
res->typed_data()[row] = ComputeRmsNorm(eps, last, &(x.typed_data()[n]), &(y->typed_data()[n]));
return ffi::Error::Success();
}
static ffi::Error RmsNormBwdImpl(ffi::Buffer<ffi::F32> res, ffi::Buffer<ffi::F32> x,
ffi::Buffer<ffi::F32> ct_y, ffi::ResultBuffer<ffi::F32> ct_x) {
auto [total, last] = GetDims(x);
if (last == 0) return ffi::Error::InvalidArgument("RmsNormBwd input must be an array");
for (int64_t n = 0, row = 0; n < total; n += last, ++row)
ComputeRmsNormBwd(last, res.typed_data()[row], &(x.typed_data()[n]), &(ct_y.typed_data()[n]),
&(ct_x->typed_data()[n]));
return ffi::Error::Success();
}
// Declare the handler signatures: attribute eps (float32), input buffers, output buffers
XLA_FFI_DEFINE_HANDLER_SYMBOL(RmsNorm, RmsNormImpl,
ffi::Ffi::Bind().Attr<float>("eps").Arg<ffi::Buffer<ffi::F32>>().Ret<ffi::Buffer<ffi::F32>>());
XLA_FFI_DEFINE_HANDLER_SYMBOL(RmsNormFwd, RmsNormFwdImpl,
ffi::Ffi::Bind().Attr<float>("eps").Arg<ffi::Buffer<ffi::F32>>()
.Ret<ffi::Buffer<ffi::F32>>().Ret<ffi::Buffer<ffi::F32>>());
XLA_FFI_DEFINE_HANDLER_SYMBOL(RmsNormBwd, RmsNormBwdImpl,
ffi::Ffi::Bind().Arg<ffi::Buffer<ffi::F32>>().Arg<ffi::Buffer<ffi::F32>>()
.Arg<ffi::Buffer<ffi::F32>>().Ret<ffi::Buffer<ffi::F32>>());
# Compile: the header path comes from jax.ffi.include_dir(); no library to link
print("FFI header directory:", short_path(jax.ffi.include_dir()))
assert (Path(jax.ffi.include_dir()) / "xla" / "ffi" / "api" / "ffi.h").exists()
so_rms = compile_so(CPP_DIR / "rms_norm.cc", CPP_DIR / "librms_norm.so", extra=["-I", jax.ffi.include_dir()])
# Load and register the three FFI targets
rms_lib = ctypes.cdll.LoadLibrary(str(so_rms))
for target, symbol in [("rms_norm", rms_lib.RmsNorm),
("rms_norm_fwd", rms_lib.RmsNormFwd),
("rms_norm_bwd", rms_lib.RmsNormBwd)]:
jax.ffi.register_ffi_target(target, jax.ffi.pycapsule(symbol), platform="cpu")
print("Registered: rms_norm, rms_norm_fwd, rms_norm_bwd")
# ffi_call: run under jit and cross-check against the jnp implementation and a NumPy float64 implementation
def rms_norm_ffi(x, eps=1e-5):
return jax.ffi.ffi_call("rms_norm", jax.ShapeDtypeStruct(x.shape, x.dtype),
vmap_method="broadcast_all")(x, eps=np.float32(eps)) # the attribute dtype must match the C++ float
def rms_norm_ref(x, eps=1e-5): # pure jnp reference implementation
return x / jnp.sqrt(jnp.mean(x ** 2, axis=-1, keepdims=True) + eps)
def rms_norm_np64(x, eps=1e-5): # NumPy float64 reference
x = np.asarray(x, np.float64)
return x / np.sqrt(np.mean(x ** 2, axis=-1, keepdims=True) + eps)
x = jax.random.normal(jax.random.key(0), (8, 64))
y_ffi = jax.jit(rms_norm_ffi)(x)
print("Max error vs. jnp:", float(jnp.max(jnp.abs(y_ffi - rms_norm_ref(x)))),
"| max error vs. NumPy f64:", float(np.max(np.abs(np.asarray(y_ffi) - rms_norm_np64(x)))))
np.testing.assert_allclose(y_ffi, rms_norm_ref(x), rtol=1e-5, atol=1e-6)
np.testing.assert_allclose(y_ffi, rms_norm_np64(x), rtol=1e-5, atol=1e-6)
np.testing.assert_allclose(rms_norm_ffi(x), y_ffi) # can also be called eagerly
hlo_txt = jax.jit(rms_norm_ffi).lower(x).as_text()
print([l.strip()[:120] for l in hlo_txt.splitlines() if "custom_call" in l][0])
assert "stablehlo.custom_call @rms_norm" in hlo_txt
# Check two common pitfalls: attribute dtype mismatch; by default ffi_call can neither be vmapped nor differentiated
try:
jax.ffi.ffi_call("rms_norm", jax.ShapeDtypeStruct(x.shape, x.dtype))(x, eps=1e-5).block_until_ready() # Python float → F64
e1 = None
except Exception as e:
e1 = e; print("eps=1e-5 (Python float) →", type(e).__name__, "|", [l for l in str(e).splitlines() if "Wrong" in l][0])
assert e1 is not None and "expected F32 but got F64" in str(e1)
plain = lambda v: jax.ffi.ffi_call("rms_norm", jax.ShapeDtypeStruct(v.shape, v.dtype))(v, eps=np.float32(1e-5))
try:
jax.vmap(plain)(x[None]); e2 = None
except NotImplementedError as e:
e2 = e; print("vmap without vmap_method →", str(e)[:90])
try:
jax.grad(lambda v: rms_norm_ffi(v).sum())(x); e3 = None
except ValueError as e:
e3 = e; print("grad →", str(e)[:110])
assert e2 is not None and e3 is not None
# ✅ custom_vjp: the forward kernel also outputs the residual s, and the backward kernel computes x̄
@partial(jax.custom_vjp, nondiff_argnums=(1,))
def rms_norm(x, eps=1e-5):
return rms_norm_ffi(x, eps)
def _rms_norm_fwd(x, eps):
y, s = jax.ffi.ffi_call(
"rms_norm_fwd",
(jax.ShapeDtypeStruct(x.shape, x.dtype), jax.ShapeDtypeStruct(x.shape[:-1], x.dtype)),
vmap_method="broadcast_all")(x, eps=np.float32(eps))
return y, (s, x) # residuals: each row's s and the input x
def _rms_norm_bwd(eps, res, ct_y):
s, x = res
ct_x = jax.ffi.ffi_call("rms_norm_bwd", jax.ShapeDtypeStruct(x.shape, x.dtype),
vmap_method="broadcast_all")(s, x, ct_y)
return (ct_x,)
rms_norm.defvjp(_rms_norm_fwd, _rms_norm_bwd)
w = jax.random.normal(jax.random.key(1), (8, 64)) # weights passed as an argument (avoids the closure-constant bug)
loss = lambda fn: (lambda x, w: jnp.sum(fn(x) * w))
g_ffi = jax.jit(jax.grad(loss(rms_norm)))(x, w)
g_ref = jax.grad(loss(rms_norm_ref))(x, w)
print("Max gradient error:", float(jnp.max(jnp.abs(g_ffi - g_ref))))
np.testing.assert_allclose(g_ffi, g_ref, rtol=1e-4, atol=1e-5)
# Finite-difference check (reverse mode, first order)
check_grads(lambda v: rms_norm(v), (x[:2, :8],), order=1, modes=["rev"], eps=1e-2, atol=1e-2, rtol=1e-2)
print("check_grads(rev) passed")
# ✅ vmap: ffi_call uses vmap_method="broadcast_all" (the kernel already treats leading dims as "rows")
xb = jax.random.normal(jax.random.key(2), (5, 8, 64))
yb = jax.vmap(rms_norm)(xb)
np.testing.assert_allclose(yb, rms_norm_np64(xb), rtol=1e-5, atol=1e-6)
# Batch dimension not in front: in_axes=1
yb1 = jax.vmap(rms_norm, in_axes=1, out_axes=1)(jnp.swapaxes(xb, 0, 1))
np.testing.assert_allclose(yb1, jnp.swapaxes(yb, 0, 1), rtol=1e-6)
# vmap ∘ grad ∘ jit: every example goes through both the fwd and bwd C++ kernels
vg = jax.jit(jax.vmap(jax.grad(loss(rms_norm)), in_axes=(0, None)))(xb, w)
vg_ref = jax.vmap(jax.grad(loss(rms_norm_ref)), in_axes=(0, None))(xb, w)
np.testing.assert_allclose(vg, vg_ref, rtol=1e-4, atol=1e-5)
print("FFI RMSNorm passes jit / grad / vmap / vmap(grad), shape:", vg.shape)
# Timing: FFI kernel vs. pure_callback (NumPy) vs. pure jnp (only the comparison is printed)
xl = jax.random.normal(jax.random.key(3), (4096, 256))
f_ffi = jax.jit(rms_norm)
f_cb = jax.jit(lambda v: jax.pure_callback(lambda a: (lambda a: a / np.sqrt(np.mean(a * a, -1, keepdims=True) + 1e-5))(np.asarray(a)),
jax.ShapeDtypeStruct(v.shape, v.dtype), v))
f_jnp = jax.jit(rms_norm_ref)
np.testing.assert_allclose(f_ffi(xl), f_jnp(xl), rtol=1e-5, atol=1e-5)
np.testing.assert_allclose(f_cb(xl), f_jnp(xl), rtol=1e-5, atol=1e-5)
t_ffi, t_cb, t_jnp = bench(f_ffi, xl), bench(f_cb, xl), bench(f_jnp, xl)
print(f"FFI (C++ scalar loop): {t_ffi:.3f} ms | pure_callback (NumPy): {t_cb:.3f} ms | pure jnp (XLA): {t_jnp:.3f} ms")
🔍 Extra¶
⚪ Optional: when you want to plug C++ / CUDA code into JAX
- FFI vs. pure_callback: pure_callback goes back to Python on every call (it's subject to the GIL and blocks fusion), and it can only run on the host. FFI produces a
custom_call; XLA calls the C++ function pointer directly without going through Python, and you can also write GPU kernels (the handler gets the CUDA stream). Use FFI when performance matters, and pure_callback for prototyping. But look at where the overhead actually goes. Measured on this machine: for a small array like(8, 64), FFI is an order of magnitude faster than pure_callback, because the fixed cost of entering and leaving Python dominates; for the large(4096, 256)array above, the two take about the same time, because NumPy is itself vectorized C code and the fixed cost gets amortized. - Naively written C++ isn't necessarily faster than XLA: in the timing above, C++ written as a scalar loop doesn't necessarily beat XLA's fused, vectorized code. FFI's real value is in reusing existing high-performance libraries, or implementing algorithms XLA can't express.
- Attribute types must match exactly:
Attr<float>requires passingnp.float32(...). A Pythonfloatis treated as F64, and the errorWrong scalar data typeonly shows up at execution time. Integer attributes are the same; usenp.int32ornp.int64. - Choosing
vmap_method: if the kernel handles arbitrary leading dimensions, use"broadcast_all"; if it can only handle a single example, use"sequential". If you don't set it, vmap fails right away. - Symbol visibility:
XLA_FFI_DEFINE_HANDLER_SYMBOLgenerates anextern "C"symbol, so ctypes can find it by name and you don't need to writeextern "C"yourself. - Turning FFI into a "primitive":
mlir.register_lowering(p, jax.ffi.ffi_lowering("rms_norm"), platform="cpu")uses an FFI target directly as a custom primitive's lowering; see Exercise 4.
Hands-on exercises¶
There are 4 exercises below. Each one has a "✍️ Your answer" cell and a "🧪 Self-check" cell. Reference solutions are in the appendix at the end of the chapter and are verified with assert. Try writing your own first, then compare with the reference solution.
Exercise 1 (linear rules for a primitive): add JVP, transpose and batching rules to custom_reduce_sum_p from 5.7, so that jax.grad, jax.vmap and jax.jit all work, and cross-check against jnp.sum. Hint: summation is linear, so the JVP is $\dot y = \mathrm{sum}(\dot x)$; the transpose just broadcasts the cotangent back to the input shape; for batching, remember that axis has to shift along with the position of the batch dimension.
Self-check interface: register three rules for custom_reduce_sum_p from 5.7, in ad.primitive_jvps, ad.primitive_transposes and batching.primitive_batchers; then call it as custom_reduce_sum(x, axis).
Exercise 2 (SciPy + pure_callback + custom_jvp): implement $\log\Gamma(x)$ with scipy.special.gammaln and pure_callback, with the derivative $\psi(x)$ from scipy.special.digamma (also via a callback). It must support vmap (pick the right vmap_method) and first-order grad, and match jax.scipy.special.gammaln and its gradient.
Self-check interface: gammaln_cb(x): defined with jax.custom_jvp, calls SciPy's $\log\Gamma$ through jax.pure_callback inside, and acts elementwise on float32 arrays of any shape.
Exercise 3 (C++ FFI elementwise kernel + gradient): write a C++ SiLU kernel $\mathrm{silu}(x) = x\,\sigma(x)$ with XLA FFI, along with its backward kernel $\bar x = \bar y\,\sigma(x)\,(1 + x(1-\sigma(x)))$. Compile and register them, hook up the gradient with custom_vjp, and cross-check against jax.nn.silu under jit, grad and vmap.
Self-check interface: silu_ffi(x): defined with jax.custom_vjp; the forward and backward passes each call one C++ kernel registered through jax.ffi; works on float32 arrays.
Exercise 4 (FFI as a primitive's lowering): create a new primitive rms_norm_prim, use jax.ffi.ffi_lowering("rms_norm") as its CPU lowering, and pass eps in as a parameter (params). Then add an eager impl (reusing the lowering) and a batching rule, and cross-check against rms_norm_np64 in eager mode, under jit and under vmap.
Self-check interface: a primitive rms_norm_prim_p, called as rms_norm_prim_p.bind(x, eps=np.float32(eps)), that applies RMSNorm over the last dimension; supports eager, jit and vmap (any in_axes).
# ✍️ Your answer: Exercise 1
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 1
def _has_batching_rule():
try:
jax.vmap(lambda v: custom_reduce_sum(v, (0,)))(jnp.ones((2, 2, 2)))
return True
except NotImplementedError: # "Batching rule for ... not implemented"
return False
_missing = [k for k, ok in [("JVP rule", custom_reduce_sum_p in ad.primitive_jvps),
("transpose rule", custom_reduce_sum_p in ad.primitive_transposes),
("batching rule", _has_batching_rule())] if not ok]
if _missing:
print(f"⏳ Exercise 1: not answered yet: custom_reduce_sum_p is missing {', '.join(_missing)}")
else:
_a = jax.random.normal(jax.random.key(21), (3, 4, 5))
_ct = jax.random.normal(jax.random.key(22), (4,))
# Linear: forward pass, JVP
np.testing.assert_allclose(custom_reduce_sum(_a, (0, 2)), np.asarray(_a).sum((0, 2)), rtol=1e-5, err_msg="The forward sum is wrong")
_, _tan = jax.jvp(lambda v: custom_reduce_sum(v, (1,)), (_a,), (jnp.ones_like(_a),))
np.testing.assert_allclose(_tan, np.full((3, 5), 4.0), rtol=1e-6, err_msg="JVP is wrong: the tangent should be sum(ẋ)")
# Transpose: broadcast the cotangent back to the input shape
_lt = jax.linear_transpose(lambda v: custom_reduce_sum(v, (0, 2)), _a)(_ct)[0]
np.testing.assert_allclose(_lt, np.broadcast_to(np.asarray(_ct)[None, :, None], (3, 4, 5)), rtol=1e-6, err_msg="Transpose rule is wrong: it should broadcast the cotangent back to the input shape")
_L = lambda v, w: jnp.sum(custom_reduce_sum(v, (0, 2)) * w)
np.testing.assert_allclose(jax.jit(jax.grad(_L))(_a, _ct), jax.grad(lambda v, w: jnp.sum(v.sum((0, 2)) * w))(_a, _ct), rtol=1e-6, err_msg="jit(grad) result is wrong")
# Batching: batch dimension in different positions
for _ax in (0, 1, 2):
np.testing.assert_allclose(jax.vmap(lambda v: custom_reduce_sum(v, (0,)), in_axes=_ax)(_a),
np.moveaxis(np.asarray(_a), _ax, 0).sum(1), rtol=1e-5,
err_msg=f"vmap(in_axes={_ax}) result is wrong: axis must shift with the batch dimension's position")
np.testing.assert_allclose(jax.vmap(jax.grad(_L), in_axes=(0, None))(jnp.stack([_a, 2 * _a]), _ct),
np.broadcast_to(np.asarray(_ct)[None, None, :, None], (2, 3, 4, 5)), rtol=1e-6, err_msg="vmap(grad) result is wrong")
check_grads(lambda v: custom_reduce_sum(v, (1,)), (_a,), order=2, modes=["fwd", "rev"], eps=1e-2)
print("✅ Exercise 1: self-check passed")
# ✍️ Your answer: Exercise 2
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 2
try:
gammaln_cb
except NameError as e:
print(f"⏳ Exercise 2: not answered yet: {e}")
else:
import scipy.special as _sps
assert isinstance(gammaln_cb, jax.custom_jvp), "Provide the derivative with custom_jvp"
_x = jnp.linspace(0.3, 9.0, 20).reshape(4, 5)
_x64 = np.asarray(_x, np.float64)
assert "pure_callback" in str(jax.make_jaxpr(gammaln_cb)(_x)), "The forward pass should call SciPy through jax.pure_callback"
np.testing.assert_allclose(jax.jit(gammaln_cb)(_x), _sps.gammaln(_x64), rtol=1e-5, atol=1e-6, err_msg="Forward values under jit are wrong")
_gfn = jax.grad(lambda v: gammaln_cb(v).sum())
np.testing.assert_allclose(jax.jit(_gfn)(_x), _sps.digamma(_x64), rtol=1e-4, atol=1e-5, err_msg="The gradient should be digamma(x)")
np.testing.assert_allclose(jax.vmap(_gfn)(_x), _sps.digamma(_x64), rtol=1e-4, atol=1e-5, err_msg="vmap(grad) result is wrong: check vmap_method")
np.testing.assert_allclose(jax.jit(jax.vmap(jax.vmap(gammaln_cb)))(_x), _sps.gammaln(_x64), rtol=1e-5, atol=1e-6, err_msg="Forward values under nested vmap are wrong")
print("✅ Exercise 2: self-check passed")
# ✍️ Your answer: Exercise 3
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 3
try:
silu_ffi
except NameError as e:
print(f"⏳ Exercise 3: not answered yet: {e}")
else:
assert isinstance(silu_ffi, jax.custom_vjp), "Hook up the gradient with custom_vjp"
_x = jax.random.normal(jax.random.key(31), (3, 7)) * 4
_w = jax.random.normal(jax.random.key(32), (3, 7))
_x64 = np.asarray(_x, np.float64)
_s = 1 / (1 + np.exp(-_x64))
_fwd_hlo = jax.jit(silu_ffi).lower(_x).as_text()
_bwd_hlo = jax.jit(jax.value_and_grad(lambda v: jnp.sum(silu_ffi(v) * _w))).lower(_x).as_text()
assert _fwd_hlo.count("stablehlo.custom_call") >= 1, "The forward pass should call an FFI kernel (no custom_call in the lowering)"
assert _bwd_hlo.count("stablehlo.custom_call") >= 2, "The backward pass should call a C++ kernel too: the value_and_grad lowering should contain two custom_calls, forward and backward"
np.testing.assert_allclose(jax.jit(silu_ffi)(_x), _x64 * _s, rtol=1e-5, atol=1e-6, err_msg="Forward values are not x·σ(x)")
_g_ref = np.asarray(_w, np.float64) * _s * (1 + _x64 * (1 - _s))
np.testing.assert_allclose(jax.jit(jax.grad(lambda v: jnp.sum(silu_ffi(v) * _w)))(_x), _g_ref, rtol=1e-4, atol=1e-5, err_msg="The gradient is wrong")
np.testing.assert_allclose(jax.vmap(jax.grad(lambda v: silu_ffi(v).sum()))(_x), _s * (1 + _x64 * (1 - _s)), rtol=1e-4, atol=1e-5, err_msg="vmap(grad) result is wrong")
print("✅ Exercise 3: self-check passed")
# ✍️ Your answer: Exercise 4
# Write your code here, then run the 🧪 Self-check below. Reference solutions are in the appendix at the end of this chapter.
# 🧪 Self-check: Exercise 4
try:
rms_norm_prim_p
except NameError as e:
print(f"⏳ Exercise 4: not answered yet: {e}")
else:
_rp = lambda v, eps=1e-5: rms_norm_prim_p.bind(v, eps=np.float32(eps))
_x = jax.random.normal(jax.random.key(41), (5, 16))
np.testing.assert_allclose(_rp(_x), rms_norm_np64(_x), rtol=1e-5, atol=1e-6, err_msg="Eager result is wrong")
np.testing.assert_allclose(jax.jit(_rp)(_x), rms_norm_np64(_x), rtol=1e-5, atol=1e-6, err_msg="jit result is wrong")
np.testing.assert_allclose(jax.jit(lambda v: _rp(v, 0.5))(_x), rms_norm_np64(_x, eps=0.5), rtol=1e-5, atol=1e-6,
err_msg="eps isn't passed to the kernel as a parameter (a different eps should change the result)")
assert "stablehlo.custom_call @rms_norm" in jax.jit(_rp).lower(_x).as_text(), "The lowering should call the FFI target rms_norm"
_x3 = jax.random.normal(jax.random.key(42), (16, 3, 6))
for _ax in (0, 1, 2):
np.testing.assert_allclose(jax.vmap(_rp, in_axes=_ax)(_x3), rms_norm_np64(np.moveaxis(np.asarray(_x3), _ax, 0)),
rtol=1e-5, atol=1e-6, err_msg=f"vmap(in_axes={_ax}) result is wrong: remember the kernel always normalizes over the last dimension")
print("✅ Exercise 4: self-check passed")
Quiz¶
Q1. In JAX 0.12, which of these hands a JAX array x to NumPy with zero copies? A. np.array(x) B. np.from_dlpack(jax.dlpack.to_dlpack(x)) C. np.from_dlpack(x) D. x.to_numpy()
Answer and explanation
C. A copies; B's to_dlpack has been removed; D doesn't exist. np.from_dlpack(x) gets the same memory through the DLPack protocol (read-only), as verified by pointer comparison in 5.2. np.asarray(x) is also zero-copy on CPU.
Q2. When does jax.dlpack.from_dlpack(np_arr) quietly copy? How do you make sure it doesn't?
Answer and explanation
When the NumPy array's start address is not 64-byte aligned (for example a slice a[1:]), the default copy=None quietly copies. Passing copy=False guarantees no copy: if zero-copy isn't possible, it raises ValueError instead.
Q3. Of the three callbacks below, which one's calls may be deleted as dead code under jit? Why? jax.debug.callback / io_callback / pure_callback
Answer and explanation
pure_callback. JAX trusts that it's a pure function, so if its output isn't used it can be deleted, and it can also be run more than once. The other two are treated as having side effects and always run. That's why you must never put logging or counting inside a pure_callback.
Q4. What happens with jax.vmap(lambda v: jax.pure_callback(f, sds, v))(xb) in 0.12? How do you fix it?
Answer and explanation
It raises NotImplementedError, because vmap_method defaults to None. Set it explicitly, for example vmap_method="expand_dims" (one call for the whole batch, when the callback supports broadcasting) or "sequential" (one call per example). The old vectorized=True no longer works, and it gets forwarded to the callback, causing a run-time error.
Q5. A custom primitive's JVP rule is written as tangent_out = 2 * x * x_dot. For jax.grad to work, must you also register a transpose rule?
Answer and explanation
No. Reverse mode transposes the linear computation in the JVP, and here the linear part only uses the built-in mul, which has its own transpose rule. You only need ad.primitive_transposes[p] when the tangent computation uses this primitive itself (for example custom_scale(x_dot, c)). Sections 5.6 and 5.9 verify the two cases respectively.
Q6. Why can't you differentiate with respect to factor in custom_scale_p.bind(x, factor=2.0)?
Answer and explanation
Keyword arguments are the primitive's static parameters (params). They're written into the jaxpr equation and take part in caching; they aren't traced values, so they have no tangents. To differentiate with respect to factor, pass it as an operand: bind(x, factor).
Q7. When calling a C++ handler with Attr<float>("eps") through XLA FFI, you write ffi_call(...)(x, eps=1e-5). What happens?
Answer and explanation
The Python float is encoded as an F64 attribute, and at execution time you get Wrong scalar data type: expected F32 but got F64. Pass np.float32(1e-5) instead, as verified in 5.10.
Q8. When hooking up C++ code, you can use ctypes + pure_callback for prototyping, but jax.ffi is now the recommended way. Name two key advantages of FFI.
Answer and explanation
① Calls don't go through Python: it compiles to a StableHLO custom_call, and XLA calls the C++ function pointer directly. There's no GIL and no round trip to the host, and it can run on GPU (the handler gets the stream). ② The handler signature is type-safe: XLA checks the dtypes of inputs, outputs and attributes, and passes shape information (dimensions()) straight to C++, so you don't have to pass sizes by hand like with ctypes.
Version notes¶
Below is a summary of the patterns and conclusions in this chapter that need attention under JAX 0.12. Each one is verified with assert at the corresponding place in the main text.
| No. | Section | Key point | Where |
|---|---|---|---|
| 1 | 5.1 | Read an array's device with the .device attribute, not .device(); use .devices() for sharded arrays | 5.1 ❌/✅ cells |
| 2 | 5.1 | Create random keys with jax.random.key(0) (typed keys), no longer jax.random.PRNGKey(0) (intro in Basics 01 Introduction, more depth in Advanced 10 Random numbers and sharding) | 5.1 ❌/✅ cells |
| 3 | 5.1 | jnp.array/asarray always copy; on CPU, jax.device_put directly aliases 64B-aligned arrays (even with may_alias=False), so modifying the source array changes the JAX array; np.asarray(jax_arr) is a zero-copy read-only view | 5.1 note + check cell |
| 4 | 5.2 | jax.dlpack.to_dlpack has been removed; use the DLPack protocol np.from_dlpack(x) / torch.from_dlpack(x) | 5.2 ❌/✅ cells |
| 5 | 5.2 | jax.dlpack.from_dlpack only accepts objects with __dlpack__; passing a raw capsule raises TypeError | 5.2 ❌ cell |
| 6 | 5.2 | x.device_buffer has been removed; get the buffer pointer with x.unsafe_buffer_pointer() / x.addressable_data(0) | 5.2 ❌/✅ cells |
| 7 | 5.2 | jax.random.normal has no device argument (passing one raises TypeError); generate first, then device_put | 5.2 ❌ cell |
| 8 | 5.2 | A zero-copy imported JAX array shares memory with the source NumPy array; modifying the source changes the JAX array's values | 5.2 note + check cell |
| 9 | 5.2 | Zero-copy NumPy→JAX requires 64-byte alignment, otherwise it quietly copies; copy=False enforces the check | 5.2 ✅ cell |
| 10 | 5.3 | jax.experimental.host_callback (id_tap/id_print/call) has been removed; use jax.debug.callback/print, io_callback, pure_callback | 5.3 note + ❌/✅ cells |
| 11 | 5.3 | On the 0.12 CPU backend all three callbacks receive jax.Array; call np.asarray before handing it to NumPy | 5.3 ✅ cells |
| 12 | 5.4 | Calling jax.grad directly through pure_callback raises ValueError: Pure callbacks do not support JVP; you need custom_jvp/custom_vjp to provide the derivative | 5.4 ❌/✅ cells |
| 13 | 5.4 | pure_callback defaults to vmap_method=None, and then vmap raises NotImplementedError; you must specify one of the 5 values | 5.4 ❌ cell + vmap_method table |
| 14 | 5.4 | See the comparison table for how the three callbacks actually behave under jit/vmap/grad/DCE (e.g. pure_callback can't be differentiated directly, and is deleted when its output is unused) | 5.3 comparison table |
| 15 | 5.5 | Primitive names are scan, cond, add (_p is only the Python variable name) | 5.5 note + cell |
| 16 | 5.5 | The compilation pipeline is jaxpr → StableHLO (MLIR) → XLA | 5.5 cell |
| 17 | 5.6 | Import Primitive from jax.extend.core (jax.core no longer exports it; ShapedArray is still in jax.core) | 5.6 ❌/✅ cells |
| 18 | 5.6 | ad.Zero.from_value doesn't exist; build zero tangents with ad.Zero(jax.typeof(y).to_tangent_aval()) | 5.6 ❌/✅ cells |
| 19 | 5.6 | ad.primitive_transposes registers a transpose rule, not a VJP; if the JVP only uses existing primitives, grad works without registering a transpose | 5.6 note + cell |
| 20 | 5.7 | A multi-output primitive must set p.multiple_results = True in addition to returning a tuple of ShapedArrays, otherwise ValueError | 5.7 note + cell |
| 21 | 5.8 | Lowering uses mlir.register_lowering + mlir.lower_fun / the StableHLO dialect (xla_client.register_translation and the XLA builder are no longer available); the GPU platform name is 'cuda', not 'gpu' | 5.8 ❌/✅ cells |
| 22 | 5.9 | Primitive.def_jvp / def_vjp, jax.core.get_aval and ad.Zero.zero_if_zero don't exist; use ad.primitive_jvps + ad.primitive_transposes | 5.9 ❌ cell |
| 23 | 5.9 | Quantities you differentiate with respect to (like custom_scale's factor) must be passed as operands bind(x, factor), not as parameters bind(x, factor=...) | 5.9 ❌/✅ cells |
| 24 | 5.9 | To support vmap a primitive needs a batching rule batching.primitive_batchers[p]; elementwise primitives can use batching.defvectorized | 5.6, 5.9 Extra |
| 25 | 5.10 | from jax.experimental import host_callback raises ImportError because the module was removed, so the whole script fails at import time | 5.10 ❌ cell |
| 26 | 5.10 | The entry point for pure_callback is jax.pure_callback (from jax.experimental import pure_callback raises ImportError) | 5.10 ❌ cell |
| 27 | 5.10 | pure_callback(..., vectorized=True) no longer works (vectorized is forwarded to the callback and causes an error); use vmap_method="expand_dims"/"broadcast_all" | 5.10 ❌/✅ cells |
| 28 | 5.10 | With x64 off, jnp.arange(..., dtype=jnp.float64) is truncated to float32 (UserWarning) | 5.10 ❌ cell |
| 29 | 5.10 | Use XLA FFI to plug in C++: jax.ffi.register_ffi_target + ffi_call / ffi_lowering (replacing xla_client.ops.CustomCall / xla.register_translation) | 5.10 FFI cells, Exercise 4 |
Chapter summary¶
- NumPy interop:
jnp.arrayalways copies;np.asarray(jax_arr)andjax.device_put(CPU, aligned arrays) may be zero-copy. Converting back to NumPy is a sync point, so do it only at the boundaries of your pipeline. - DLPack: 0.12 only has the protocol-style interface:
np.from_dlpack(x),jax.dlpack.from_dlpack(obj, copy=...), and the arrays' own__dlpack__/__dlpack_device__. Zero-copy needs the same device and 64-byte alignment. Shared memory means changes to the source array "leak through" into the JAX array. - The callback trio replaces
host_callback:debug.callbackfor observing,io_callbackfor side effects (ordered=Trueguarantees order), andpure_callbackfor pure computation (you must setvmap_method, and it may be removed by DCE). None of them supports autodiff on its own; add derivatives withcustom_jvp/custom_vjp. - Primitives are the smallest unit of every transformation. A custom primitive needs
def_impl,def_abstract_eval(returning aShapedArray),mlir.register_lowering,ad.primitive_jvps, plusad.primitive_transposesandbatching.primitive_batchersas needed. All of these are semi-internal APIs. - C++ integration: use ctypes +
pure_callbackfor prototyping; in production use XLA FFI: C++ handler → compile againstjax.ffi.include_dir()→pycapsule+register_ffi_target→ffi_call(vmap_method=...)→custom_vjpfor gradients.
Cheat sheet¶
| API (JAX 0.12) | Purpose |
|---|---|
x.device / x.devices() | The array's device (attribute / set) |
jax.device_put(np_arr, device) | Put on a given device (may alias NumPy memory on CPU) |
np.asarray(x) / np.array(x) | JAX → NumPy: zero-copy read-only view / copy |
jax.device_get(tree) | Pull a whole pytree back to the host |
np.from_dlpack(x) | Zero-copy JAX → NumPy (DLPack protocol) |
jax.dlpack.from_dlpack(obj, device=None, copy=None) | External array → JAX; copy=False enforces zero-copy |
x.__dlpack__() / x.__dlpack_device__() | Producer side of the DLPack protocol |
x.unsafe_buffer_pointer() | Address of the underlying buffer (for checking zero-copy) |
jax.debug.print(fmt, **kw) / jax.debug.callback(fn, *args) | Debug printing / observing intermediate values (no return value) |
jax.experimental.io_callback(fn, result_sds, *args, ordered=) | Host callback with side effects |
jax.pure_callback(fn, result_sds, *args, vmap_method=) | Pure host callback; vmap_method ∈ {sequential, sequential_unrolled, expand_dims, broadcast_all, legacy_vectorized} |
jax.effects_barrier() | Wait until all callbacks have run |
jax.ShapeDtypeStruct(shape, dtype) | Describe the output of a callback / FFI call |
jax.extend.core.Primitive(name) | Create a primitive (semi-internal) |
p.def_impl / p.def_abstract_eval | Eager implementation / abstract evaluation (returns jax.core.ShapedArray) |
p.multiple_results = True | Multi-output primitive |
mlir.register_lowering(p, rule, platform=) | Register a lowering (jax.interpreters.mlir) |
mlir.lower_fun(fn, multiple_results=) | Generate a lowering from a jnp implementation |
jax.extend.mlir.dialects.stablehlo | Hand-write StableHLO ops |
ad.primitive_jvps[p] / ad.primitive_transposes[p] | JVP / transpose rules (jax.interpreters.ad) |
ad.Zero(aval) / ad.is_undefined_primal(x) | Symbolic zero tangent / test for a linear input in a transpose |
batching.primitive_batchers[p] / batching.defvectorized(p) | vmap rule / one-line batching for elementwise primitives |
xla.apply_primitive | Implement the eager impl via the lowering |
jax.eval_shape(fn, *sds) | Abstract evaluation only (for debugging shapes) |
jax.test_util.check_grads(f, args, order, modes) | Numerically check differentiation rules |
jax.ffi.include_dir() | Directory of the XLA FFI C++ headers |
jax.ffi.pycapsule(fn_ptr) / jax.ffi.register_ffi_target(name, cap, platform=) | Wrap and register a C++ handler |
jax.ffi.ffi_call(name, result_sds, vmap_method=)(*args, **attrs) | Call an FFI target (attribute dtypes must match) |
jax.ffi.ffi_lowering(name) | Use an FFI target as a primitive's lowering |
XLA_FFI_DEFINE_HANDLER_SYMBOL / ffi::Ffi::Bind().Attr<>().Arg<>().Ret<>() | Declare the handler signature on the C++ side |
Appendix: reference solutions¶
Write your own answer in the "✍️ Your answer" cell in the main text first and check it with the "🧪 Self-check"; come here to compare once the self-check passes or you're truly stuck. Running the solutions below in order overwrites any functions with the same names you defined in the answer cells.
Exercise 1¶
See "Hands-on exercises" in the main text for the task.
# Reference solution: Exercise 1 — JVP / transpose / batching for custom_reduce_sum
def _reduce_sum_jvp(primals, tangents, *, axis):
x, = primals; t, = tangents
y = custom_reduce_sum_p.bind(x, axis=axis)
if type(t) is ad.Zero:
return y, ad.Zero(jax.typeof(y).to_tangent_aval())
return y, custom_reduce_sum_p.bind(t, axis=axis) # linear: the tangent uses this primitive too
def _reduce_sum_transpose(ct, x, *, axis):
assert ad.is_undefined_primal(x)
shape = x.aval.shape
if type(ct) is ad.Zero:
return (ad.Zero(x.aval),)
return (jnp.broadcast_to(jnp.expand_dims(ct, axis), shape),) # broadcast back to the input shape
def _reduce_sum_batch(args, dims, *, axis):
x, = args; bd, = dims
x = jnp.moveaxis(x, bd, 0) # move the batch dim to 0
new_axis = tuple(a + 1 for a in axis) # all original axes shift right by 1
return custom_reduce_sum_p.bind(x, axis=new_axis), 0
ad.primitive_jvps[custom_reduce_sum_p] = _reduce_sum_jvp
ad.primitive_transposes[custom_reduce_sum_p] = _reduce_sum_transpose
batching.primitive_batchers[custom_reduce_sum_p] = _reduce_sum_batch
a = jax.random.normal(jax.random.key(0), (2, 3, 4))
wt = jax.random.normal(jax.random.key(1), (3,))
L1 = lambda a, wt: jnp.sum(custom_reduce_sum(a, (0, 2)) * wt)
L1_ref = lambda a, wt: jnp.sum(jnp.sum(a, (0, 2)) * wt)
np.testing.assert_allclose(jax.jit(jax.grad(L1))(a, wt), jax.grad(L1_ref)(a, wt), rtol=1e-6)
np.testing.assert_allclose(jax.vmap(lambda v: custom_reduce_sum(v, (1,)), in_axes=2)(a), jnp.moveaxis(a.sum(1), 1, 0), rtol=1e-6)
np.testing.assert_allclose(jax.vmap(jax.grad(L1), in_axes=(0, None))(a[None].repeat(3, 0), wt),
jax.vmap(jax.grad(L1_ref), in_axes=(0, None))(a[None].repeat(3, 0), wt), rtol=1e-6)
check_grads(lambda v: custom_reduce_sum(v, (1,)), (a,), order=2, modes=["fwd", "rev"], eps=1e-2)
print("Exercise 1 passed")
Exercise 2¶
See "Hands-on exercises" in the main text for the task.
# Reference solution: Exercise 2 — SciPy gammaln / digamma through pure_callback, derivative via custom_jvp
import scipy.special as sps
def _scipy_cb(fn, x):
return jax.pure_callback(lambda a: fn(np.asarray(a)).astype(a.dtype),
jax.ShapeDtypeStruct(x.shape, x.dtype), x,
vmap_method="expand_dims") # SciPy ufuncs handle any shape
@jax.custom_jvp
def gammaln_cb(x):
return _scipy_cb(sps.gammaln, x)
@gammaln_cb.defjvp
def _gammaln_jvp(p, t):
(x,), (xd,) = p, t
return gammaln_cb(x), _scipy_cb(sps.digamma, x) * xd
xs = jnp.linspace(0.5, 6.0, 12)
np.testing.assert_allclose(jax.jit(gammaln_cb)(xs), jax.scipy.special.gammaln(xs), rtol=1e-5)
np.testing.assert_allclose(jax.grad(lambda v: gammaln_cb(v).sum())(xs),
jax.grad(lambda v: jax.scipy.special.gammaln(v).sum())(xs), rtol=1e-4)
xs2 = xs.reshape(3, 4)
np.testing.assert_allclose(jax.vmap(jax.grad(lambda v: gammaln_cb(v).sum()))(xs2),
sps.digamma(np.asarray(xs2, np.float64)), rtol=1e-4)
print("Exercise 2 passed")
Exercise 3¶
See "Hands-on exercises" in the main text for the task.
# Reference solution: Exercise 3 (C++ part) — SiLU forward / backward FFI kernels, written to cpp/silu.cc
(CPP_DIR / "silu.cc").write_text(r'''// Exercise 3: SiLU forward / backward FFI kernels (elementwise)
#include <cmath>
#include <cstdint>
#include "xla/ffi/api/ffi.h"
namespace ffi = xla::ffi;
static inline float Sigmoid(float v) { return 1.0f / (1.0f + std::exp(-v)); }
static ffi::Error SiluImpl(ffi::Buffer<ffi::F32> x, ffi::ResultBuffer<ffi::F32> y) {
const int64_t n = x.element_count();
const float* xp = x.typed_data();
float* yp = y->typed_data();
for (int64_t i = 0; i < n; ++i) yp[i] = xp[i] * Sigmoid(xp[i]);
return ffi::Error::Success();
}
// ct_x = ct_y * s * (1 + x * (1 - s)), s = sigmoid(x)
static ffi::Error SiluBwdImpl(ffi::Buffer<ffi::F32> x, ffi::Buffer<ffi::F32> ct_y,
ffi::ResultBuffer<ffi::F32> ct_x) {
const int64_t n = x.element_count();
const float* xp = x.typed_data();
const float* gp = ct_y.typed_data();
float* op = ct_x->typed_data();
for (int64_t i = 0; i < n; ++i) {
float s = Sigmoid(xp[i]);
op[i] = gp[i] * s * (1.0f + xp[i] * (1.0f - s));
}
return ffi::Error::Success();
}
XLA_FFI_DEFINE_HANDLER_SYMBOL(Silu, SiluImpl,
ffi::Ffi::Bind().Arg<ffi::Buffer<ffi::F32>>().Ret<ffi::Buffer<ffi::F32>>());
XLA_FFI_DEFINE_HANDLER_SYMBOL(SiluBwd, SiluBwdImpl,
ffi::Ffi::Bind().Arg<ffi::Buffer<ffi::F32>>().Arg<ffi::Buffer<ffi::F32>>().Ret<ffi::Buffer<ffi::F32>>());
''')
print("Wrote", short_path(CPP_DIR / "silu.cc"))
# Reference solution: Exercise 3 — compile and register the SiLU kernels, custom_vjp + vmap
so_silu = compile_so(CPP_DIR / "silu.cc", CPP_DIR / "libsilu.so", extra=["-I", jax.ffi.include_dir()])
silu_lib = ctypes.cdll.LoadLibrary(str(so_silu))
jax.ffi.register_ffi_target("silu", jax.ffi.pycapsule(silu_lib.Silu), platform="cpu")
jax.ffi.register_ffi_target("silu_bwd", jax.ffi.pycapsule(silu_lib.SiluBwd), platform="cpu")
@jax.custom_vjp
def silu_ffi(x):
return jax.ffi.ffi_call("silu", jax.ShapeDtypeStruct(x.shape, x.dtype), vmap_method="broadcast_all")(x)
silu_ffi.defvjp(
lambda x: (silu_ffi(x), x), # the residual is just x
lambda x, g: (jax.ffi.ffi_call("silu_bwd", jax.ShapeDtypeStruct(x.shape, x.dtype),
vmap_method="broadcast_all")(x, g),))
xv = jax.random.normal(jax.random.key(5), (6, 10)) * 3
np.testing.assert_allclose(jax.jit(silu_ffi)(xv), jax.nn.silu(xv), rtol=1e-5, atol=1e-6)
gv = jax.jit(jax.grad(lambda v, w: jnp.sum(silu_ffi(v) * w)))(xv, xv)
np.testing.assert_allclose(gv, jax.grad(lambda v, w: jnp.sum(jax.nn.silu(v) * w))(xv, xv), rtol=1e-4, atol=1e-5)
np.testing.assert_allclose(jax.vmap(jax.grad(lambda v: silu_ffi(v).sum()))(xv),
jax.vmap(jax.grad(lambda v: jax.nn.silu(v).sum()))(xv), rtol=1e-4, atol=1e-6)
print("Exercise 3 passed")
Exercise 4¶
See "Hands-on exercises" in the main text for the task.
# Reference solution: Exercise 4 — an FFI target as a custom primitive's lowering
from jax.interpreters import xla as _xla
rms_norm_prim_p = Primitive("rms_norm_prim")
rms_norm_prim_p.def_abstract_eval(lambda x, *, eps: ShapedArray(x.shape, x.dtype))
mlir.register_lowering(rms_norm_prim_p, jax.ffi.ffi_lowering("rms_norm"), platform="cpu") # params become FFI attributes
rms_norm_prim_p.def_impl(partial(_xla.apply_primitive, rms_norm_prim_p)) # eager: jit this one primitive on the fly
def _rms_prim_batch(args, dims, *, eps):
x, = args; bd, = dims
# The kernel treats every dim except the last as "rows": the batch dim just must not be last, so move it to 0
return rms_norm_prim_p.bind(jnp.moveaxis(x, bd, 0), eps=eps), 0
batching.primitive_batchers[rms_norm_prim_p] = _rms_prim_batch
rms_prim = lambda v: rms_norm_prim_p.bind(v, eps=np.float32(1e-5))
xq = jax.random.normal(jax.random.key(7), (4, 32))
np.testing.assert_allclose(rms_prim(xq), rms_norm_np64(xq), rtol=1e-5, atol=1e-6) # eager
np.testing.assert_allclose(jax.jit(rms_prim)(xq), rms_norm_np64(xq), rtol=1e-5, atol=1e-6) # jit
xq3 = jax.random.normal(jax.random.key(8), (32, 3, 4))
np.testing.assert_allclose(jax.vmap(rms_prim, in_axes=2)(xq3),
rms_norm_np64(np.moveaxis(np.asarray(xq3), 2, 0)), rtol=1e-5, atol=1e-6)
print(jax.make_jaxpr(rms_prim)(xq))
print("Exercise 4 passed")