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

  1. Know which conversions between jax.Array and numpy.ndarray copy, which share memory, and where the sync points are.
  2. 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.
  3. Learn the three replacements for the removed host_callback: jax.debug.callback, io_callback and pure_callback, and how each behaves under jit/vmap/grad.
  4. 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.
  5. Use XLA FFI (jax.ffi) to plug a real compiled C++ function into JAX, then add gradient and vmap support.

Sections

SectionTopicOriginal slug
5.1Integrating JAX with NumPyintegrating-jax-numpy
5.2Zero-copy data sharing with DLPackzero-copy-dlpack
5.3host_callback → three kinds of callbackscalling-external-code-host-callback
5.4Side-effect-free calls with jax.pure_callbackusing-jax-pure-callback
5.5Introduction to JAX primitivesintroduction-jax-primitives
5.6Defining custom primitivesdefining-custom-primitives
5.7Implementing abstract evaluation rulesimplementing-abstract-evaluation
5.8Implementing lowering rules for backendsimplementing-lowering-rules
5.9Differentiation rules for custom primitives (+ batching)differentiation-rules-custom-primitives
5.10Practice: 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.

  • 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_method and 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
Python 3.14.7 | JAX 0.12.0.1 | devices [CpuDevice(id=0)]
C++ directory: cpp

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 as jit, grad and vmap, 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

  1. Convert as little as possible in inner loops, and never inside an @jax.jit function.
  2. Convert only at the boundaries of your computation: bring data in with jnp.array / jax.device_put at the start, and take it out with np.array at the end.
  3. On multiple devices, use jax.device_put to state placement explicitly.
  4. 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.
  5. Prefer jax.numpy in JAX code. Passing NumPy arrays straight into jnp functions 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]
Original NumPy type: <class 'numpy.ndarray'> (implicitly on the host CPU)
Converted type: <class 'jaxlib._jax.ArrayImpl'> | isinstance(jax.Array): True
Device: cpu: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)
TypeError: 'jaxlib._jax.Device' object is not callable
# ✅ 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
cpu:0 {CpuDevice(id=0)}
Available devices: [CpuDevice(id=0)]
Device after explicit placement: cpu:0
# ❌ 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)
Legacy key: [0 0] | dtype: uint32 | shape: (2,)
# ✅ 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)
Typed key: Array((), dtype=key<fry>) overlaying:
[0 0] | dtype: key<fry>
JAX array: [4.245284   5.0505295  0.13281113] | device: cpu:0
NumPy type: <class 'numpy.ndarray'> | value: [4.245284   5.0505295  0.13281113]

⚠️ Note: jnp.array and jax.device_put have 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 passing may_alias=False doesn'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; only np.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
jnp.array                    shares memory with NumPy: False
jnp.asarray                  shares memory with NumPy: False
jax.device_put               shares memory with NumPy: True
device_put(may_alias=False)  shares memory with NumPy: True
After modifying the source: jnp.array[0] = 1.0 | device_put[0] = 42.0
np.asarray zero-copy: True | writeable: False
np.array   zero-copy: False | writeable: True
# 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))
convert every step: 25.04 ms | convert only at the end: 21.57 ms | ratio 1.16x
Reading: syncing every step is slower, as expected.

🔍 Extra

🟢 Must-read

  • You can't apply NumPy functions to tracers inside jit: np.sin(tracer) raises TracerArrayConversionError. 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 a np_arr.copy() first, or use jnp.array instead.
  • jax.device_get(tree) pulls a whole pytree back to the host in one call, which is easier than calling np.array on 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())
TracerArrayConversionError : The numpy.ndarray conversion method __array__() was called on traced array with shape float32[3]

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 PyCapsule in Python) that holds: ① a pointer to the underlying buffer; ② metadata: device type and ID (CPU/CUDA/ROCm…), dtype, number of dimensions ndim, 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

  1. 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.
  2. 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.
  3. 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.
  4. Mutability: JAX arrays are immutable. If the consumer modifies the underlying buffer in place, it breaks that assumption.
  5. Stream ordering (advanced): on a GPU, if the two sides use different CUDA streams, calling block_until_ready() or torch.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 to torch.from_dlpack or jax.dlpack.from_dlpack, with the warning that "a capsule can only be consumed once". In JAX 0.12, jax.dlpack.to_dlpack has been removed; the module only has from_dlpack and is_supported_dtype. jax.dlpack.from_dlpack also no longer accepts raw capsules; passing one raises TypeError.

The current way is the DLPack protocol: hand the array object directly to the consumer, for example np.from_dlpack(x) or torch.from_dlpack(x), and the consumer calls x.__dlpack__() and x.__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; use x.unsafe_buffer_pointer() or x.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)")
jax.dlpack.to_dlpack(x)                       → AttributeError: module 'jax.dlpack' has no attribute 'to_dlpack'
x.device_buffer.unsafe_buffer_pointer()       → AttributeError: arr.device_buffer has been deprecated. Use arr.addressable_data(0)
jax.dlpack.from_dlpack(x.__dlpack__())        → TypeError: The array passed to from_dlpack must have __dlpack__ and __dlpack_device__ methods.
jax.random.normal(key, shape, device=...)     → TypeError: normal() got an unexpected keyword argument 'device'
Public members of jax.dlpack: ['from_dlpack', 'is_supported_dtype']
(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)
__dlpack_device__: (<DLDeviceType.kDLCPU: 1>, 0)
__dlpack__() returns: PyCapsule
JAX pointer  : 537202262016
NumPy pointer: 537202262016 | writeable: False
# ✅ 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
from_dlpack shares memory: True | jnp.from_dlpack shares memory: True
unaligned array, from_dlpack shares memory: False
copy=False → Specified input which requires a copy since the source data buffer is not aligned to 64 bytes.  However copy=F

⚠️ 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
before: j[0] = 1.0 | after modifying the source: j[0] = 99.0
After deleting the original JAX array, NumPy 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")
jnp.array (copy): 1.577 ms | from_dlpack (zero-copy): 0.0645 ms | 24x 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) and jax.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 passing copy=False then raises. Put the other way round: copy=False guarantees 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 (returns arg unchanged) while triggering fn on the host, for side effects like printing and logging. id_print is 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, so result_shape_dtypes is 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_spec returns None, and importing it raises ImportError. 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_transform has 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 receive jax.Array (ArrayImpl). So inside a callback, do np.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
find_spec: None
ImportError: cannot import name 'host_callback' from 'jax.experimental' (<venv>/lib/python3.14/site-packages/jax/experimental/__init__.py)
# ✅ 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)
[host callback] shape=(8, 8) first 5=[0.939  1.2573 1.156  1.0641 0.8711]
Final result: 0.5013
# ✅ 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])
intermediate y = [5. 6. 7.]
# ✅ 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)
Final result: [0.75726515 0.64076287 0.29125583 0.64076287 0.75726515]
# ✅ 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
Write order: [0, 1, 2, 3, 4] | 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?

CallbackjitvmapgradWhen 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 supportedKept
io_callback(ordered=True)✅❌ ValueError: can't vmap an ordered callback❌ ValueErrorKept
pure_callback (no vmap_method)✅❌ NotImplementedError❌ ValueError: suggests custom_jvpDeleted, never runs
pure_callback(vmap_method="expand_dims")✅✅ (one call that receives the whole batch)❌ ValueErrorDeleted
# 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
callback                    jit   vmap                  vmap calls    grad        runs when unused
debug.callback              OK    OK                    4             OK          1
io_callback                 OK    OK                    4             ValueError  1
io_callback(ordered)        OK    ValueError            0             ValueError  1
pure_callback(default)      OK    NotImplementedError   0             ValueError  0
pure_callback(expand_dims)  OK    OK                    1             ValueError  0

🔍 Extra

🟢 Must-read

  • How to choose: to just peek at an intermediate value, use jax.debug.print or jax.debug.callback. To write files, send metrics, or rely on execution order, use io_callback(ordered=True). For pure computation that needs a return value, use pure_callback. pure_callback is 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.callback only sees primal values: it doesn't take part in differentiation and sees neither tangents nor cotangents. To print a gradient, call jax.debug.print on the result of jax.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_jvp or jax.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.grad should work directly". In JAX 0.12 that isn't the case:

  1. vmap: vmap_method defaults to None, and then vmap raises NotImplementedError right away, so you must set it explicitly. The valid values in 0.12 are "sequential", "sequential_unrolled", "expand_dims", "broadcast_all" and "legacy_vectorized".
  2. 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 with jax.custom_jvp or jax.custom_vjp.
  3. The old argument vectorized=True is no longer a pure_callback parameter in 0.12. It gets forwarded unchanged to the callback as part of **kwargs, and you get TypeError: ... 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)
Output:
 [[ 4.6199546   4.8222537  -0.27386507]
 [ 0.7643099   1.5264554  -1.6242192 ]
 [-0.44588774  2.4433482   2.8974433 ]]
Batched output shape: (10, 3, 3)
# ❌ 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])
vmap → NotImplementedError: vmap is only supported for the pure_callback primitive when vmap_method is one of 'sequential', 'sequential_unrolled', 'expand_dims', 'broadcast_all',
grad → ValueError: Pure callbacks do not support JVP. Please use `jax.custom_jvp` to use callbacks while taking gradients.
# ✅ 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)
Gradients:
 [[0.8963548 0.1220305 2.8149228]
 [2.9938226 2.9690719 2.1271498]
 [2.7596536 2.7605276 2.5746362]]
# 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
Multiple outputs: res1=Array(20., dtype=float32), res2=Array(6., dtype=float32)

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_methodHow many times the callback runsShape the callback receivesWhen 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"OnceBatched arguments are $(B, *s)$; unbatched arguments get a leading dimension of size 1The callback supports NumPy broadcasting (the most common choice)
"broadcast_all"OnceAll arguments are broadcast to $(B, *s)$The callback needs every input to have the same batch dimension
"legacy_vectorized"OnceBatched 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])
sequential           called 4 time(s), (a, b) shapes received = ((3,), (3,))  output (4, 3)
sequential_unrolled  called 4 time(s), (a, b) shapes received = ((3,), (3,))  output (4, 3)
expand_dims          called 1 time(s), (a, b) shapes received = ((4, 3), (1, 3))  output (4, 3)
broadcast_all        called 1 time(s), (a, b) shapes received = ((4, 3), (4, 3))  output (4, 3)
legacy_vectorized    called 1 time(s), (a, b) shapes received = ((4, 3), (3,))  output (4, 3)
Invalid value → vmap_method must be on of the allowed methods ['sequential', 'sequential_unrolled', 'expand_dims', 'broadcast_all', 'leg

🔍 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 sequential calls 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 error at 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_jvp is also computed with a pure_callback, first-order grad works, but jax.hessian runs 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))
hessian(v1) → ValueError: Pure callbacks do not support JVP. Please use `jax.custom_jvp` to use callbacks 
Hessian diagonal after the fix: [-0.         -0.32719472 -0.6183698  -0.841471  ]
# 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
Wrong return shape → JaxRuntimeError
Callback runs when the output is unused: 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")
pure_callback: 2.030 ms | native jnp: 1.529 ms | ratio 1.3x

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):

  1. Transformations act on primitives: jit traces a function into a jaxpr and hands it to XLA to compile; grad relies on each primitive's differentiation rule; vmap relies on each primitive's batching rule. Transformations never look at Python code, only at primitives.
  2. 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.
  3. 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 scan and cond, not scan_p and cond_p; the primitive registered by jnp.add is also called add, not add_p. *_p is 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")
{ lambda ; a:f32[3] b:f32[3]. let
    c:f32[3] = sin a
    d:f32[3] = mul c b
    e:f32[] = dot_general[
      dimension_numbers=(([0], [0]), ([], []))
      preferred_element_type=float32
    ] a b
    f:f32[3] = add d e
    g:f32[] = reduce_sum[axes=(0,) out_sharding=None] f
  in (g,) }
Primitive sequence: ['sin', 'mul', 'dot_general', 'add', 'reduce_sum']
Control-flow primitives: {'cond', 'scan'}
P.scan_p.name = scan | P.cond_p.name = cond | jax.lax.add_p.name = 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
sin has a JVP rule       : True
mul has a transpose rule : True
primitive_batchers doesn't support `in`: argument of type 'PrimitiveBatchersProxy' is not a container or iterable
sin has a batching rule  : True
module @jit_f attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} {
  func.func public @main(%arg0: tensor<3xf32>, %arg1: tensor<3xf32>) -> (tensor<f32> {jax.result_info = "result"}) {
    %0 = stablehlo.sine %arg0 : tensor<3xf32>
    %1 = stablehlo.multiply %0, %arg1 : tensor<3xf32>
    %2 = stablehlo.dot_general %arg0, %arg1, contracting_dims = [0] x [0], precision = [DEFAULT, DEFAULT] : (tensor<3xf32>, tensor<3xf32>) -> tensor<f32>
    %3 = stablehlo.broadcast_in_dim %2, dims = [] : (tensor<f32>) -> tensor<3xf32>
    %4 = stablehlo.add %1, %3 : tensor<3xf32>
    %cst = stablehlo.constant dense<0.000000e+00> : tensor<f32>
    %5 = stablehlo.reduce(%4 init: %cst) applies stablehlo.add across dimensions = [0] : (tensor<3xf32>, tensor<f32>) -> tensor<f32>
    return %5 : tensor<f32>
  }
}

🔍 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):
RuleWhere it's registered (0.12)Transformation it servesError when missing
impl (eager implementation)p.def_implDirect calls (not under jit)Evaluation rule ... not implemented
Abstract evaluationp.def_abstract_evalAll tracing: jit / eval_shape / grad / vmapAbstract evaluation ... not implemented
loweringmlir.register_loweringjit (compilation)MLIR translation rule ... not found
JVPad.primitive_jvps[p]jvp / gradDifferentiation rule ... not implemented
Transposead.primitive_transposes[p]grad / vjp (only linear primitives need it)Transpose rule ... not implemented
batchingbatching.primitive_batchers[p]vmapBatching rule ... not implemented
  • In 0.12, batching.primitive_batchers is a write-only proxy: batching.primitive_batchers[p] = rule wraps the rule with a new signature and stores it in batching.fancy_primitive_batchers. It doesn't support in and can't be read; query the latter instead. Also, when none of the arguments is batched, the wrapper doesn't call your rule but calls bind directly. 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_callback or jax.ffi.ffi_call plus custom_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:

  1. Create the primitive object; its name shows up in the jaxpr.
  2. Write a wrapper function for users that calls p.bind(*args, **params). bind doesn'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 as some_parameter) must be passed to bind as keyword arguments.
  3. Register the abstract evaluation rule with def_abstract_eval (5.7).
  4. Register the backend lowering rule with mlir.register_lowering (5.8).
  5. 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 Primitive raises ImportError, because jax.core no longer exports Primitive. The recommended path is from jax.extend.core import Primitive. ShapedArray is still at jax.core.ShapedArray; it's not in jax.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"))
ImportError: cannot import name 'Primitive' from 'jax.core' (<venv>/lib/python3.14/site-packages/jax/core.py)
jax.core has Primitive: False | has ShapedArray: True
jax.extend.core has Primitive: True | has ShapedArray: False
# ✅ 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]}")
eager: [ 2.  5. 10.]
{ lambda ; a:f32[3]. let b:f32[3] = my_custom_op[some_parameter=1.0] a in (b,) }
jit  → NotImplementedError: MLIR translation rule for primitive 'my_custom_op' not found for platform cpu
grad → NotImplementedError: Differentiation rule for 'my_custom_op' not implemented
vmap → NotImplementedError: Batching rule for 'my_custom_op' not implemented
# ✅ 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])
jit: [ 2.  5. 10.]
# ❌ 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)
ad.Zero has from_value: False
AttributeError: type object 'Zero' has no attribute 'from_value'
# ✅ 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")
grad: [2. 4. 6.] | jit(grad): [2. 4. 6.]
vmap / vmap(grad) passed

⚠️ Note: ad.primitive_transposes registers 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 as 2 * x * x_dot (the built-in mul), and mul comes with its own transpose rule, so registering only the JVP is enough for jax.grad to 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:

  1. Static analysis: shape mismatches and type errors are caught before anything is compiled or run.
  2. Transformation compatibility: vmap relies on it to derive batched shapes, and grad relies on it to build the structure of the backward pass.
  3. 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
eval_shape: ShapeDtypeStruct(shape=(4, 5), dtype=float32) | impl calls: 0
TypeError at tracing time: custom_add_one only supports floating-point inputs, got int32
# ✅ 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))
ShapeDtypeStruct(shape=(3, 4), dtype=float32)
{ lambda ; a:f32[2,3,4]. let
    b:f32[2,3] = custom_reduce_sum[axis=(2,)] a
  in (b,) }

⚠️ Note: for a primitive with multiple outputs, having abstract evaluation return a tuple of ShapedArrays isn't enough; you must also set p.multiple_results = True first. Otherwise, returning a tuple from abstract evaluation raises ValueError: ... 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)
ValueError: minmax_bad.abstract_eval() method should return a tuple or a list iff minmax_bad.multiple_results.
min, max = -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, like float(x) or np.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. The custom_reduce_sum above first normalizes axis into a sorted tuple, so equivalent ways of writing it hit the same cache entry.
  • jax.eval_shape is 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_translation and the XLA builder (builder.Add, builder.GetShape and so on) all belong to the old XLA translation-rule API. In JAX 0.12, importing jax.lib.xla_client fails, and jax.interpreters.xla has no register_translation (only a few members like apply_primitive remain).

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 is rule(ctx, *mlir_values, **params) -> list[ir.Value], and it produces StableHLO. The simplest way is mlir.lower_fun(<jnp implementation>, multiple_results=...). For fine control, write ops by hand with jax.extend.mlir.dialects.stablehlo. To call external C++, use jax.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")
ImportError: cannot import name 'xla_client' from 'jax.lib' (<venv>/lib/python3.14/site-packages/jax/lib/__init__.py)
Public members of jax.interpreters.xla: ['Backend', 'apply_primitive', 'canonicalize_dtype_handlers']
Has register_translation: False
# ✅ 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
module @jit_bind attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} {
  func.func public @main(%arg0: tensor<3xf32>) -> (tensor<3xf32> {jax.result_info = "result"}) {
    %cst = stablehlo.constant dense<0.000000e+00> : tensor<f32>
    %0 = stablehlo.broadcast_in_dim %cst, dims = [] : (tensor<f32>) -> tensor<3xf32>
    %1 = stablehlo.compare GT, %arg0, %0, FLOAT : (tensor<3xf32>, tensor<3xf32>) -> tensor<3xi1>
    %cst_0 = stablehlo.constant dense<0.000000e+00> : tensor<f32>
    %2 = stablehlo.broadcast_in_dim %cst_0, dims = [] : (tensor<f32>) -> tensor<3xf32>
    %3 = call @_where(%1, %arg0, %2) : (tensor<3xi1>, tensor<3xf32>, tensor<3xf32>) -> tensor<3xf32>
    return %3 : tensor<3xf32>
  }
  func.func private @_where(%arg0: tensor<3xi1>, %arg1: tensor<3xf32>, %arg2: tensor<3xf32>) -> tensor<3xf32> {
    %0 = stablehlo.select %arg0, %arg1, %arg2 : tensor<3xi1>, tensor<3xf32>
    return %0 : tensor<3xf32>
  }
}

# ✅ 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])
module @jit_bind attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} {
  func.func public @main(%arg0: tensor<4xf32>) -> (tensor<4xf32> {jax.result_info = "result"}) {
    %cst = stablehlo.constant dense<0.000000e+00> : tensor<f32>
    %0 = stablehlo.broadcast_in_dim %cst, dims = [] : (tensor<f32>) -> tensor<4xf32>
    %1 = stablehlo.maximum %0, %arg0 : tensor<4xf32>
    return %1 : tensor<4xf32>
  }
}

🔍 Extra

⚪ Optional: when you hand-write a lowering, or register different implementations per platform

  • Values of platform: None means 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, calls jax.core.get_aval inside abstract evaluation, and builds zeros with ad.Zero.zero_if_zero. In JAX 0.12, Primitive has no def_jvp / def_vjp methods at all, and neither jax.core.get_aval nor ad.Zero.zero_if_zero exists. 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 to bind as 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])
Primitive has def_jvp: False | has def_vjp: False
jax.core has get_aval: False | ad.Zero has zero_if_zero: False
jax.ad has Zero: False (jax.ad is now a different, new namespace)
def_jvp → AttributeError: 'Primitive' object has no attribute 'def_jvp'
bind(x, factor=...) → TypeError: <lambda>() got an unexpected keyword argument 'factor'
# ✅ 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])
JVP: [2. 4. 6.] [0.7 1.4 2.1]
dL/dfactor (no transpose rule yet): 56.0
dL/dx → Transpose rule (for reverse-mode differentiation) for 'custom_scale' not implemented
# ✅ 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.")
dL/dx = [ 8. 16. 24.] | dL/dfactor = 56.0
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))
custom_scale passes jit / jvp / grad / vmap / vmap(grad) / check_grads
{ lambda ; a:f32[3] b:f32[]. let
    c:f32[3] = custom_scale a b
    d:f32[3] = mul c c
    _:f32[] = reduce_sum[axes=(0,) out_sharding=None] d
    e:f32[3] = broadcast_in_dim 1.0:f32[]
    f:f32[3] = mul c e
    g:f32[3] = mul e c
    h:f32[3] = add_any g f
    i:f32[3] = custom_scale h b
  in (i,) }

🔍 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 as check_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 with ad.is_undefined_primal, and their .aval gives the shape. Nonlinear inputs are ordinary values. The return value has one entry per operand, with None in the positions of nonlinear inputs.
  • Zero tangents: ad.Zero(aval) is a symbolic zero and takes no memory. In a rule, check type(t) is ad.Zero first, and don't do arithmetic on it directly. When you need a real array of zeros, use ad.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_vjp is enough (5.4, 5.10), and you never need to touch ad.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:

  1. C++ implementation: use extern "C" to prevent C++ name mangling, so ctypes can find the symbol by name; compute input[i]*input[i] + 10.0 elementwise over a double array.
  2. Compile to a shared library: g++ -shared -fPIC -o custom_op.so custom_op.cpp on Linux, g++ -shared -o custom_op.dylib custom_op.cpp on macOS, and the matching MinGW or MSVC command on Windows.
  3. ctypes wrapper: load the library with ctypes.CDLL and set argtypes/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].
  4. pure_callback integration: 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.
  5. The alternative (outline only): define a Primitive → abstract evaluation → per-backend lowering (mentioning xla_client.ops.CustomCall and xla.register_translation) → JVP/transpose → user function calling bind.
  6. 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:

  1. from jax.experimental import host_callback as hcb raises ImportError right away. Even if it's never used after the import, the whole script fails at import time;
  2. from jax.experimental import pure_callback raises ImportError; the correct entry point is jax.pure_callback;
  3. pure_callback(..., vectorized=True): vectorized is no longer a pure_callback parameter. It gets forwarded to the callback as a kwarg and fails at run time. Use vmap_method="expand_dims" or "broadcast_all" instead;
  4. 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.CustomCall and xla.register_translation (and the earlier popular approach of exporting a PyCapsule with pybind11) are the old way of hooking in custom calls. 0.12 recommends XLA FFI: on the C++ side, define a type-safe handler with XLA_FFI_DEFINE_HANDLER_SYMBOL from xla/ffi/api/ffi.h; on the Python side, register it with jax.ffi.pycapsule + jax.ffi.register_ffi_target, then call it with jax.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;
    }
}
}
Writing cpp/custom_op.cpp
# 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()
Apple clang version 21.0.0 (clang-2100.4.1.1)
$ /usr/bin/clang++ -std=c++17 -O2 -shared -fPIC -o cpp/libcustom_op.so cpp/custom_op.cpp
Compiled → libcustom_op.so (16816 bytes)
# 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)
NumPy wrapper test: [11. 14. 19.]
# ❌ 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
from jax.experimental import host_callback as hcb → ImportError
from jax.experimental import pure_callback → ImportError
Requested float64, actual dtype = float32 | warning: Explicitly requested dtype float64 requested in arange is not availabl
vectorized=True → JaxRuntimeError | root cause: custom_op_numpy() got an unexpected keyword argument 'vectorized'
# ✅ 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))
eager: [11. 14. 19. 26.] | jit: [11. 14. 19. 26.]
grad: [2. 4. 6. 8.]

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:

  1. C++ side: write an ordinary function ffi::Error Impl(attributes..., ffi::Buffer<T> inputs..., ffi::ResultBuffer<T> outputs...), then declare its signature with XLA_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++.
  2. Compile: the headers live under jax.ffi.include_dir(). You only need the headers; no library to link.
  3. Python side: load with ctypes.cdll.LoadLibrary → wrap the function pointer with jax.ffi.pycapsule(lib.<symbol>) → register with jax.ffi.register_ffi_target(<name>, <capsule>, platform="cpu").
  4. Call: jax.ffi.ffi_call(<name>, <output ShapeDtypeStruct>, vmap_method=...)(*inputs, **attributes). It produces a StableHLO custom_call, so it works with jit out of the box, and there's no round trip through Python like with callbacks.
  5. Gradients: ffi_call isn't differentiable, so use jax.custom_vjp with 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>>());
Writing cpp/rms_norm.cc
# 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 header directory: <venv>/lib/python3.14/site-packages/jaxlib/include
$ /usr/bin/clang++ -std=c++17 -O2 -shared -fPIC -I <venv>/lib/python3.14/site-packages/jaxlib/include -o cpp/librms_norm.so cpp/rms_norm.cc
Compiled → librms_norm.so (93664 bytes)
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
Max error vs. jnp: 4.76837158203125e-07 | max error vs. NumPy f64: 3.514383175584612e-07
%0 = stablehlo.custom_call @rms_norm(%arg0) {backend_config = "", mhlo.backend_config = {eps = 9.99999974E-6 : f32}, ope
# 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
eps=1e-5 (Python float) → JaxRuntimeError | Wrong scalar data type: expected F32 but got F64
vmap without vmap_method → vmap is only supported for the ffi_call primitive when vmap_method is one of 'sequential',
grad → The FFI call to `rms_norm` cannot be differentiated. You can use `jax.custom_jvp` or `jax.custom_vjp` to add s
# ✅ 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")
Max gradient error: 4.76837158203125e-07
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)
FFI RMSNorm passes jit / grad / vmap / vmap(grad), shape: (5, 8, 64)
# 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")
FFI (C++ scalar loop): 0.740 ms | pure_callback (NumPy): 0.704 ms | pure jnp (XLA): 0.087 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 passing np.float32(...). A Python float is treated as F64, and the error Wrong scalar data type only shows up at execution time. Integer attributes are the same; use np.int32 or np.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_SYMBOL generates an extern "C" symbol, so ctypes can find it by name and you don't need to write extern "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")
⏳ Exercise 1: not answered yet: custom_reduce_sum_p is missing JVP rule, transpose rule, batching rule
# ✍️ 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")
⏳ Exercise 2: not answered yet: name 'gammaln_cb' is not defined
# ✍️ 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")
⏳ Exercise 3: not answered yet: name 'silu_ffi' is not defined
# ✍️ 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")
⏳ Exercise 4: not answered yet: name 'rms_norm_prim_p' is not defined

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.SectionKey pointWhere
15.1Read an array's device with the .device attribute, not .device(); use .devices() for sharded arrays5.1 ❌/✅ cells
25.1Create 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
35.1jnp.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 view5.1 note + check cell
45.2jax.dlpack.to_dlpack has been removed; use the DLPack protocol np.from_dlpack(x) / torch.from_dlpack(x)5.2 ❌/✅ cells
55.2jax.dlpack.from_dlpack only accepts objects with __dlpack__; passing a raw capsule raises TypeError5.2 ❌ cell
65.2x.device_buffer has been removed; get the buffer pointer with x.unsafe_buffer_pointer() / x.addressable_data(0)5.2 ❌/✅ cells
75.2jax.random.normal has no device argument (passing one raises TypeError); generate first, then device_put5.2 ❌ cell
85.2A zero-copy imported JAX array shares memory with the source NumPy array; modifying the source changes the JAX array's values5.2 note + check cell
95.2Zero-copy NumPy→JAX requires 64-byte alignment, otherwise it quietly copies; copy=False enforces the check5.2 ✅ cell
105.3jax.experimental.host_callback (id_tap/id_print/call) has been removed; use jax.debug.callback/print, io_callback, pure_callback5.3 note + ❌/✅ cells
115.3On the 0.12 CPU backend all three callbacks receive jax.Array; call np.asarray before handing it to NumPy5.3 ✅ cells
125.4Calling jax.grad directly through pure_callback raises ValueError: Pure callbacks do not support JVP; you need custom_jvp/custom_vjp to provide the derivative5.4 ❌/✅ cells
135.4pure_callback defaults to vmap_method=None, and then vmap raises NotImplementedError; you must specify one of the 5 values5.4 ❌ cell + vmap_method table
145.4See 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
155.5Primitive names are scan, cond, add (_p is only the Python variable name)5.5 note + cell
165.5The compilation pipeline is jaxpr → StableHLO (MLIR) → XLA5.5 cell
175.6Import Primitive from jax.extend.core (jax.core no longer exports it; ShapedArray is still in jax.core)5.6 ❌/✅ cells
185.6ad.Zero.from_value doesn't exist; build zero tangents with ad.Zero(jax.typeof(y).to_tangent_aval())5.6 ❌/✅ cells
195.6ad.primitive_transposes registers a transpose rule, not a VJP; if the JVP only uses existing primitives, grad works without registering a transpose5.6 note + cell
205.7A multi-output primitive must set p.multiple_results = True in addition to returning a tuple of ShapedArrays, otherwise ValueError5.7 note + cell
215.8Lowering 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
225.9Primitive.def_jvp / def_vjp, jax.core.get_aval and ad.Zero.zero_if_zero don't exist; use ad.primitive_jvps + ad.primitive_transposes5.9 ❌ cell
235.9Quantities 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
245.9To support vmap a primitive needs a batching rule batching.primitive_batchers[p]; elementwise primitives can use batching.defvectorized5.6, 5.9 Extra
255.10from jax.experimental import host_callback raises ImportError because the module was removed, so the whole script fails at import time5.10 ❌ cell
265.10The entry point for pure_callback is jax.pure_callback (from jax.experimental import pure_callback raises ImportError)5.10 ❌ cell
275.10pure_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
285.10With x64 off, jnp.arange(..., dtype=jnp.float64) is truncated to float32 (UserWarning)5.10 ❌ cell
295.10Use 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

  1. NumPy interop: jnp.array always copies; np.asarray(jax_arr) and jax.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.
  2. 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.
  3. The callback trio replaces host_callback: debug.callback for observing, io_callback for side effects (ordered=True guarantees order), and pure_callback for pure computation (you must set vmap_method, and it may be removed by DCE). None of them supports autodiff on its own; add derivatives with custom_jvp/custom_vjp.
  4. Primitives are the smallest unit of every transformation. A custom primitive needs def_impl, def_abstract_eval (returning a ShapedArray), mlir.register_lowering, ad.primitive_jvps, plus ad.primitive_transposes and batching.primitive_batchers as needed. All of these are semi-internal APIs.
  5. C++ integration: use ctypes + pure_callback for prototyping; in production use XLA FFI: C++ handler → compile against jax.ffi.include_dir() → pycapsule + register_ffi_target → ffi_call(vmap_method=...) → custom_vjp for 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_evalEager implementation / abstract evaluation (returns jax.core.ShapedArray)
p.multiple_results = TrueMulti-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.stablehloHand-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_primitiveImplement 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 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 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"))
Wrote cpp/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")
$ /usr/bin/clang++ -std=c++17 -O2 -shared -fPIC -I <venv>/lib/python3.14/site-packages/jaxlib/include -o cpp/libsilu.so cpp/silu.cc
Compiled → libsilu.so (70824 bytes)
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")
{ lambda ; a:f32[4,32]. let b:f32[4,32] = rms_norm_prim[eps=1e-05] a in (b,) }
Exercise 4 passed