第 5 章 互操作性与自定义操作

进阶课程「JAX 高级编程:性能、优化与规模化」学习笔记 · 第 5 章 原课程链接:https://apxml.com/zh/courses/advanced-jax(本章:chapter-5-jax-interoperability-custom-operations) 运行环境:Python 3.14 · JAX / jaxlib 0.12.0.1(仅 CPU) · NumPy · SciPy · Apple clang++

阅读说明:这是 API 变化最多的一章。host_callback、to_dlpack、xla_client、register_translation 和 def_jvp/def_vjp 在 JAX 0.12 里都已经用不了。正文代码全部按 0.12 的写法来写。过时写法放在 ❌ 单元里,后面紧跟 ✅ 推荐写法,每条结论都用 assert 验证过。

⚠️ 半内部 API 提醒:5.5~5.9 用到的 jax.extend.core.Primitive、jax.interpreters.{mlir, ad, batching} 和 jax.core.ShapedArray 都是半内部 API,JAX 不保证它们跨版本稳定。这里的导入路径只保证在 0.12.0.1 上能用。升级 JAX 之后,先把相关单元跑一遍再往下看。

学习目标

  1. 弄清 jax.Array 与 numpy.ndarray 互转时,哪些操作复制、哪些共享内存、哪里是同步点。
  2. 用 DLPack 协议(__dlpack__ / np.from_dlpack / jax.dlpack.from_dlpack)实现零拷贝,并用指针比较证明内存确实共享。
  3. 掌握 host_callback 被删除之后的三种替代:jax.debug.callback、io_callback、pure_callback,以及它们在 jit/vmap/grad 下的不同行为。
  4. 理解原语(Primitive)是 JAX 变换的最小单元,能从零定义一个原语,并依次补上 impl、抽象求值、lowering、JVP、转置和 batching 规则,让它通过 jit、grad 和 vmap。
  5. 用 XLA FFI(jax.ffi) 把真正编译出来的 C++ 函数接进 JAX,并补上梯度和 vmap 支持。

小节目录

小节主题原文 slug
5.1JAX 与 NumPy 的集成integrating-jax-numpy
5.2使用 DLPack 实现零拷贝数据共享zero-copy-dlpack
5.3host_callback → 三种回调calling-external-code-host-callback
5.4使用 jax.pure_callback 进行无副作用调用using-jax-pure-callback
5.5JAX 原语简介introduction-jax-primitives
5.6定义自定义原语defining-custom-primitives
5.7实现抽象求值规则implementing-abstract-evaluation
5.8为后端实现转换(lowering)规则implementing-lowering-rules
5.9制定自定义原语的求导规则(+ batching)differentiation-rules-custom-primitives
5.10实践:整合 C++ 函数(ctypes 回调 → XLA FFI)practice-integrating-cpp-function

之后依次是:动手练习、测验、本章更新说明、本章小结、速查表。

与入门课的关联

⚠️ 环境注意:jaxlib 0.12.0.1 的 CPU 后端有一个已知 bug:jit 闭包捕获的较大常量矩阵在 matmul(W @ X)之后做 reduce,会悄悄算错(复现与规避见 入门 00 目录与导读 的「本环境已知问题」,根因分析见 进阶 02 §2.3)。所以本章的权重和数据都作为函数参数传入,关键数值用 NumPy float64 对拍。

# 公共准备:导入、版本检查、C++ 目录、计时与对齐分配工具
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__, "| 设备", jax.devices())
assert jax.__version__.startswith("0.12")
assert jax.default_backend() == "cpu", "本章的示例按 CPU 编写(在 Colab 上请使用 CPU 运行时)"

NB_DIR = Path.cwd()                 # notebook 所在的目录
CPP_DIR = NB_DIR / "cpp"            # C++ 源码与 .so 输出目录
CPP_DIR.mkdir(exist_ok=True)
def short_path(s):
    """打印时把本机绝对路径缩短(notebook 目录 → 相对路径,虚拟环境 → <venv>),输出里不出现个人目录"""
    return str(s).replace(str(NB_DIR) + "/", "").replace(sys.prefix, "<venv>")

print("C++ 目录:", short_path(CPP_DIR))

def bench(f, *args, n=20):
    """平均耗时(毫秒):先预热一次,再循环 n 次,最后 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):
    """分配起始地址按 align 字节对齐的一维 NumPy 数组(DLPack 零拷贝导入 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):
    """JAX 数组 j 与 NumPy 数组 n 是否指向同一块内存(比较起始指针)"""
    return j.unsafe_buffer_pointer() == n.ctypes.data
Python 3.14.7 | JAX 0.12.0.1 | 设备 [CpuDevice(id=0)]
C++ 目录: cpp

5.1 JAX 与 NumPy 的集成

5.1 JAX 与 NumPy 的集成

jax.numpy 特意模仿了 NumPy 的 API,但这两种数组本质上不一样,互相转换时要留意性能。

两种数组的区别

  • numpy.ndarray:数据放在主机 CPU 内存里,运算由 CPU 直接执行(eager)。
  • jax.Array:数据由 JAX 运行时管理。它可以放在 CPU 上,但真正的优势是能放在 GPU/TPU 这类加速器上。jit、grad、vmap 等变换和编译后的计算,处理的都是它。
  • 两者放的位置不同、用途也不同,所以转换时要多留心。

NumPy → JAX

  • 最直接的办法是 jnp.array(np_arr):JAX 从主机内存读出数据,放到默认设备上(有 GPU/TPU 就放那里,没有就放 CPU)。
  • 想明确指定放到哪个设备,用 jax.device_put(np_arr, device=...)。分布式计算、或者需要保证数据局部性时会用到。

JAX → NumPy

  • np.array(jax_arr) 把结果拉回主机,用于保存、绘图或交给其他库。
  • 性能上要注意:这个转换分两步。① 等设备把这个数组算完。② 把数据从设备内存拷回主机内存。JAX 默认是异步调度的(入门课 01),所以这一步是一个同步点:Python 会一直等到数据就绪。在性能敏感的循环里频繁转换,会把速度拖慢很多。

课程给出的 5 条最佳实践

  1. 内层循环里尽量少转换,@jax.jit 函数内部更是不要转换。
  2. 只在计算流程的边界转换:入口用 jnp.array / jax.device_put 转入,出口用 np.array 转出。
  3. 多设备场景下用 jax.device_put 明确表达放置意图。
  4. np.array(jax_array) 会悄悄等计算完成。如果只想同步、不想拷数据(比如精确计时),直接调用 .block_until_ready()。
  5. JAX 代码里优先用 jax.numpy。直接把 NumPy 数组传给 jnp 函数也能跑,但其中的隐式转换可能把性能开销藏起来。

总之,要把转换(尤其是 JAX→NumPy)当成可能很贵的操作,设计程序时有意识地安排它们放在哪里。

# ✅ 推荐写法:NumPy → JAX(对应课程代码块 1;设备用 .device 属性)
numpy_arr = np.array([1.0, 2.0, 3.0], dtype=np.float32)
print("原始 NumPy 类型:", type(numpy_arr), "(隐含在主机 CPU 上)")

jax_arr = jnp.array(numpy_arr)
print("转换后类型:", type(jax_arr), "| isinstance(jax.Array):", isinstance(jax_arr, jax.Array))
print("所在设备:", jax_arr.device)

assert isinstance(jax_arr, jax.Array) and jax_arr.dtype == jnp.float32
assert jax_arr.device == jax.devices()[0]
原始 NumPy 类型: <class 'numpy.ndarray'> (隐含在主机 CPU 上)
转换后类型: <class 'jaxlib._jax.ArrayImpl'> | isinstance(jax.Array): True
所在设备: cpu:0
# ❌ 旧写法(过时/错误):jax_arr.device() —— 0.12 中 .device 是属性(Device 对象),不能调用
try:
    print(jax_arr.device())
except TypeError as e:
    print("TypeError:", e)
TypeError: 'jaxlib._jax.Device' object is not callable
# ✅ 推荐写法:.device 属性(单设备数组);.devices() 返回设备集合(分片数组也适用)
print(jax_arr.device, jax_arr.devices())
assert jax_arr.devices() == {jax_arr.device}

# 课程代码块 2:用 jax.device_put 显式放置
available_devices = jax.devices()
print("可用设备:", 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("显式放置后的设备:", jax_arr_explicit.device)
assert jax_arr_explicit.device == target_device
cpu:0 {CpuDevice(id=0)}
可用设备: [CpuDevice(id=0)]
显式放置后的设备: cpu:0
# ❌ 旧写法(过时/错误):jax.random.PRNGKey(0) 返回 uint32[2] 的“旧式原始 key”,0.12 推荐类型化 key
try:
    old_key = jax.random.PRNGKey(0)
    print("旧式 key:", old_key, "| dtype:", old_key.dtype, "| shape:", old_key.shape)
except Exception as e:
    print(type(e).__name__, e)
旧式 key: [0 0] | dtype: uint32 | shape: (2,)
# ✅ 推荐写法:jax.random.key 生成类型化 key;JAX → NumPy(课程代码块 3)
key = jax.random.key(0)
print("类型化 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 数组:", jax_arr, "| 设备:", jax_arr.device)
numpy_result = np.array(jax_arr)     # 同步点:等待计算完成,并复制回主机
print("NumPy 类型:", type(numpy_result), "| 值:", numpy_result)
assert isinstance(numpy_result, np.ndarray)
np.testing.assert_array_equal(numpy_result, jax_arr)
类型化 key: Array((), dtype=key<fry>) overlaying:
[0 0] | dtype: key<fry>
JAX 数组: [4.245284   5.0505295  0.13281113] | 设备: cpu:0
NumPy 类型: <class 'numpy.ndarray'> | 值: [4.245284   5.0505295  0.13281113]

⚠️ 注意:jnp.array 和 jax.device_put 的内存语义并不一样,区别不只是「一个用默认设备、一个显式指定设备」。在 JAX 0.12 的 CPU 后端上:

  • jnp.array(np_arr) / jnp.asarray(np_arr):总是复制。
  • jax.device_put(np_arr):只要源数组按 64 字节对齐,它就直接别名 NumPy 的内存,传 may_alias=False 也挡不住。本机 macOS 上几 KB 以上的 NumPy 数组都是 64 字节对齐的,很小的数组不一定,具体取决于平台的内存分配器。这时再去改 NumPy 源数组,这个“不可变”的 JAX 数组会跟着变。
  • 反过来,np.asarray(jax_arr) 是零拷贝的只读视图,np.array(jax_arr) 才会复制。
# 动手验证:各种转换到底复制不复制(用指针比较判断)
big = aligned_zeros(1 << 20); big[:] = 1.0           # 4 MB,64 字节对齐
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} 与 NumPy 共享内存: {shares_memory(arr, big)}")

big[0] = 42.0                                           # 修改 NumPy 源数组
print("修改源数组后: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   # ⚠️ “不可变”的 JAX 数组被改了
assert shares_memory(a_put_noalias, big)                        # 0.12.0.1 实测:may_alias=False 也别名

# 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 零拷贝:", n_view.ctypes.data == y.unsafe_buffer_pointer(), "| 可写:", n_view.flags.writeable)
print("np.array   零拷贝:", n_copy.ctypes.data == y.unsafe_buffer_pointer(), "| 可写:", 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                    与 NumPy 共享内存: False
jnp.asarray                  与 NumPy 共享内存: False
jax.device_put               与 NumPy 共享内存: True
device_put(may_alias=False)  与 NumPy 共享内存: True
修改源数组后:jnp.array[0] = 1.0 | device_put[0] = 42.0
np.asarray 零拷贝: True | 可写: False
np.array   零拷贝: False | 可写: True
# 同步点的代价:每步都转回 NumPy,对比只在最后转一次(计时只打印,不做 assert)
@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])        # 每步都同步并复制
    return x

def loop_convert_at_end(x, n=100):
    for _ in range(n):
        x = step(x)                           # 异步排队,不阻塞
    return x

t_each = bench(loop_convert_each_step, x0, n=3)
t_end = bench(loop_convert_at_end, x0, n=3)
print(f"每步转换: {t_each:.2f} ms | 只在末尾转换: {t_end:.2f} ms | 比值 {t_each / t_end:.2f}x")
print("解读:" + ("每步同步更慢,符合预期。" if t_each > t_end * 1.1 else
      "本机仅 CPU:设备内存就是主机内存,“拷回”几乎免费,异步排队的收益也有限,两者差距不明显,甚至可能反转。"
      "同步点的代价主要出现在 GPU/TPU 上,那里每次都要等待设备并走 PCIe 传输。"))
np.testing.assert_allclose(loop_convert_each_step(x0), loop_convert_at_end(x0))
每步转换: 23.06 ms | 只在末尾转换: 15.02 ms | 比值 1.54x
解读:每步同步更慢,符合预期。

🔍 补充

🟢 必读

  • jit 内部不能用 NumPy 函数处理 tracer:np.sin(tracer) 会抛出 TracerArrayConversionError。NumPy 要的是具体数值,可追踪的时候只有抽象值(见 02_JIT编译)。真要在 jit 里调 NumPy/SciPy,就得用 5.3、5.4 讲的回调。
  • 零拷贝视图有利也有弊:np.asarray(jax_arr) 不复制,但只读。jax.device_put(np_arr) 在 CPU 上可能别名,传进去之后如果还要原地改这个 NumPy 数组,先 np_arr.copy() 一份,或者改用 jnp.array。
  • jax.device_get(tree) 可以一次把整个 pytree 拉回主机,比逐个叶子调用 np.array 更省事。
# 验证:在 jit 里对 tracer 调 NumPy 函数
@jax.jit
def bad(x):
    return np.sin(x)      # 追踪时 x 是 tracer,NumPy 无法把它转成具体数组

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 使用 DLPack 实现零拷贝数据共享

5.2 使用 DLPack 实现零拷贝数据共享

问题背景:JAX 和 NumPy、PyTorch、TensorFlow、CuPy 之间交换数据,常见做法是复制,比如 JAX → 主机 NumPy → PyTorch → 再搬回 GPU。如果是放在加速器上的大数组,这样来回搬、重复占内存,代价很大。就算 JAX 和 PyTorch 在同一块 GPU 上,在两套内存之间复制几个 GB 的数据,也既费时间,又占宝贵的内存带宽。

DLPack 标准:一种和编程语言无关的“内存中张量”的描述规范。数据在同一个设备上时,支持它的库之间可以零拷贝交换数据,真正传过去的只是指针和元数据。

  • 导出方生成一个胶囊(Python 里是 PyCapsule),里面有:① 指向底层缓冲区的指针;② 元数据,也就是设备类型和 ID(CPU/CUDA/ROCm…)、dtype、维数 ndim、形状、步幅(strides)和字节偏移。
  • 接收方读出指针和元数据,直接包装这块现成的内存,做成自己的张量对象,数据本身不复制。

课程列出的 5 条注意事项

  1. 设备兼容:双方用的是同一个物理设备,才能零拷贝。DLPack 没法直接在 CPU 和 GPU 之间共享,数据得先搬到目标设备上。
  2. 数据同步:JAX 是异步执行的,导出前先 block_until_ready()。从别的框架(比如 CUDA 上的 PyTorch)导入前,也要确认对方的计算流已经同步,不然可能读到还没写完的数据。
  3. 生命周期与所有权:内存一般归原来的框架所有。导入的对象还在用的时候,原张量必须活着,不然会出现悬空指针。
  4. 可变性:JAX 数组是不可变的。如果接收方原地改了底层缓冲区,JAX 的这个前提就被打破了。
  5. 流顺序(高级):GPU 上双方用的是不同的 CUDA stream 时,光调用 block_until_ready() 或 torch.cuda.synchronize() 可能还不能保证对方看得到数据,需要用流事件来协调。

⚠️ 注意:旧资料里常见的写法是先用 jax.dlpack.to_dlpack(x) 导出胶囊,再把胶囊交给 torch.from_dlpack 或 jax.dlpack.from_dlpack,还要强调“胶囊只能用一次”。在 JAX 0.12 里,jax.dlpack.to_dlpack 已经移除,模块里只剩 from_dlpack 和 is_supported_dtype。jax.dlpack.from_dlpack 也不再接受裸胶囊,传进去会报 TypeError。

现在的写法是 DLPack 协议:直接把数组对象交给消费方,比如 np.from_dlpack(x) 或 torch.from_dlpack(x),由消费方自己去调 x.__dlpack__() 和 x.__dlpack_device__()。用户碰不到胶囊,“只能用一次”的问题也就没有了。另外,x.device_buffer.unsafe_buffer_pointer() 也已移除,改用 x.unsafe_buffer_pointer() 或 x.addressable_data(0).unsafe_buffer_pointer()。

# ❌ 旧写法(过时/错误):to_dlpack / device_buffer / from_dlpack(胶囊) / random.normal(device=) 在 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} → 仍可用")
    except Exception as e:
        print(f"{desc:45s} → {type(e).__name__}: {str(e).splitlines()[0][:90]}")
print("jax.dlpack 公开成员:", [a for a in dir(jax.dlpack) if not a.startswith("_")])
print("(本环境没有安装 torch / cupy,这两段示例在这里跑不了)")
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'
jax.dlpack 公开成员: ['from_dlpack', 'is_supported_dtype']
(本环境没有安装 torch / cupy,这两段示例在这里跑不了)
# ✅ 推荐写法:DLPack 协议 —— JAX → NumPy 零拷贝
x = jax.random.normal(jax.random.key(0), (1024, 1024))
x.block_until_ready()                       # 导出前先同步(注意事项 2)
print("__dlpack_device__:", x.__dlpack_device__())
print("__dlpack__() 返回:", type(x.__dlpack__()).__name__)

n = np.from_dlpack(x)                       # NumPy 自己调用 x.__dlpack__()
print("JAX 指针  :", x.unsafe_buffer_pointer())
print("NumPy 指针:", n.ctypes.data, "| 可写:", 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                # 导出的是只读视图,保护 JAX 的不可变性
np.testing.assert_array_equal(n, x)
__dlpack_device__: (<DLDeviceType.kDLCPU: 1>, 0)
__dlpack__() 返回: PyCapsule
JAX 指针  : 492067356672
NumPy 指针: 492067356672 | 可写: False
# ✅ 推荐写法: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)                   # 等价入口
print("from_dlpack 共享内存:", shares_memory(j1, src), "| jnp.from_dlpack 共享内存:", shares_memory(j2, src))
assert shares_memory(j1, src) and shares_memory(j2, src)

# 对齐要求:起始地址不是 64 字节对齐时,默认(copy=None)会静默复制;copy=False 直接报错
unaligned = aligned_zeros(1025)[1:]         # 偏移 4 字节
assert unaligned.ctypes.data % 64 != 0
j_u = jax.dlpack.from_dlpack(unaligned)
print("未对齐数组 from_dlpack 共享内存:", 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 共享内存: True | jnp.from_dlpack 共享内存: True
未对齐数组 from_dlpack 共享内存: False
copy=False → Specified input which requires a copy since the source data buffer is not aligned to 64 bytes.  However copy=F

⚠️ 注意:零拷贝导入得到的 JAX 数组,并不是真的“改不了”。通过 JAX 的 API 确实改不了它,但它和源 NumPy 数组共享同一块内存:直接改源数组,JAX 数组的值也会跟着变。这会破坏 JAX 的函数式假设,比如 jit 缓存的常量,或者已经排队的异步计算读到的值。所以零拷贝导入之后,就把源数组当成“冻结”的,别再动它。

# 验证:零拷贝导入后修改源数组 → JAX 数组随之改变;生命周期由消费方持有引用保证
src = aligned_zeros(1024); src[:] = 1.0
j = jax.dlpack.from_dlpack(src)
before = float(j[0])
src[0] = 99.0
print("修改前 j[0] =", before, "| 修改源数组后 j[0] =", float(j[0]))
assert before == 1.0 and float(j[0]) == 99.0

# 生命周期:删掉 JAX 数组后,NumPy 视图仍然有效(DLPack 的 deleter 机制让消费方持有生产方的引用)
x = jnp.full((1000,), 3.0)
view = np.from_dlpack(x)
del x
import gc; gc.collect()
print("原 JAX 数组删除后,NumPy 视图求和 =", view.sum())
assert view.sum() == 3000.0
修改前 j[0] = 1.0 | 修改源数组后 j[0] = 99.0
原 JAX 数组删除后,NumPy 视图求和 = 3000.0
# 计时:零拷贝 vs 复制(64 MB;只打印比较结果)
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(复制): {t_copy:.3f} ms | from_dlpack(零拷贝): {t_dl:.4f} ms | 快 {t_copy / t_dl:.0f}x")
jnp.array(复制): 1.556 ms | from_dlpack(零拷贝): 0.0695 ms | 快 22x

🔍 补充

⚪ 选读:要和 PyTorch、CuPy 等框架交换数据时

  • 和 PyTorch 互转的现代写法(本环境没有 torch,这一条没有实际跑过):torch.from_dlpack(jax_arr) 和 jax.dlpack.from_dlpack(torch_tensor) 都直接传对象,不需要中间胶囊。在 GPU 上,消费方会通过 __dlpack__(stream=...) 协商流同步。
  • 支持的 dtype:用 jax.dlpack.is_supported_dtype(dtype) 查询。
  • 跨设备:jax.dlpack.from_dlpack(obj, device=..., copy=...) 可以指定目标设备。目标设备和源设备不同时一定会复制,这时传 copy=False 会报错。反过来说,用 copy=False 就能确保“真的是零拷贝”;不加的话,零拷贝可能悄悄变成了复制。

5.3 调用外部 CPU/GPU 代码:从 host_callback 到三种回调

5.3 使用 jax.experimental.host_callback 调用外部 CPU/GPU 代码

XLA 编译让 JAX 跑得快。但有时我们需要在计算执行期间临时跳出编译环境,到主机 CPU 上跑一段任意的 Python 代码,比如调试、记录中间值、和外部硬件交互,或者调用没有 JAX 实现的库。课程介绍的是 jax.experimental.host_callback,它主要提供两个函数:

  • id_tap(fn, arg, tap_with_transform=...):在计算图里相当于一个恒等函数(原样返回 arg),同时触发主机端的 fn,用来做打印、记日志这类副作用。id_print 是专门用来打印的简化版。
  • call(fn, args, result_shape_dtypes=...):调用主机函数,并把返回值送回设备,接回 JAX 的数据流。追踪的时候就得知道输出的形状和 dtype,所以必须提供 result_shape_dtypes。

运行机制(发生在执行的时候,不是追踪的时候):① 执行到回调指令时,设备暂停;② 参数从设备拷到主机;③ 在主机上运行 Python 函数;④(只有 call 才有)把返回值拷回设备;⑤ 设备继续执行。这样走一个来回,就带来了同步点和传输开销。

课程列出的注意事项:① 性能开销大;② 对 XLA 来说是个看不透的优化屏障,回调两边的算子没法融合;③ 回调也是异步调度的,打印可能要等到需要结果时才出现;④ 允许副作用,函数就不纯了;⑤ 默认不可微;⑥ 和变换的配合:jit 能用,vmap 下行为复杂,pmap 下每台主机都会执行一次。

什么时候用:当作最后的手段,或者用在对性能不敏感的任务上,比如调试、日志、对接遗留代码、和硬件交互。其他选择有:用 jax.numpy/jax.lax 重写;纯函数式的外部调用用 jax.pure_callback(5.4);高性能的 C++/CUDA 用自定义原语或 FFI(5.6~5.10)。

⚠️ 注意:jax.experimental.host_callback(id_tap、id_print、call)在 JAX 0.12 里已经删除了。importlib.util.find_spec 返回 None,导入直接报 ImportError。它的功能拆给了三个正式 API: | 旧 API | 0.12 替代 | 适用场景 | |---|---|---| | id_tap / id_print | jax.debug.callback / jax.debug.print | 调试打印,不返回值 | | call(有副作用) | jax.experimental.io_callback(..., ordered=True/False) | 有副作用、需要返回值或保证顺序 | | call(纯函数) | jax.pure_callback(..., vmap_method=...) | 纯计算,可被 vmap,可被 DCE |

tap_with_transform 没有对应的替代。另外,很容易以为主机函数收到的是 NumPy 数组,但在 0.12 的 CPU 后端上,三种回调收到的其实都是 jax.Array(ArrayImpl)。所以回调里要先 np.asarray(arg),再交给 NumPy/SciPy。

# ❌ 旧写法(过时/错误):from jax.experimental import host_callback —— 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)
# ✅ 推荐写法:id_tap → jax.debug.callback(对应课程代码块 1)
records = []
def log_intermediate_data(arg):
    records.append(type(arg))
    arg = np.asarray(arg)                    # 回调收到 jax.Array,先转 NumPy
    print(f"[主机回调] 形状={arg.shape} 前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)   # 不返回值,不影响数据流
    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()                        # 等所有回调执行完,避免打印漏到别的单元
print("最终结果:", 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)
[主机回调] 形状=(8, 8) 前5个=[0.939  1.2573 1.156  1.0641 0.8711]
最终结果: 0.5013
# ✅ 推荐写法:id_print → jax.debug.print(格式字符串,懒格式化)
@jax.jit
def simple_print_example(x):
    y = x + 5.0
    jax.debug.print("中间值 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])
中间值 y = [5. 6. 7.]
# ✅ 推荐写法:host_callback.call(纯计算)→ jax.pure_callback(对应课程代码块 2)
def external_cpu_calculation(data, parameter):
    data = np.asarray(data)                  # 主机端:普通 NumPy 计算
    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),   # 输出形状与 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)
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)
最终结果: [0.75726515 0.64076287 0.29125583 0.64076287 0.75726515]
# ✅ 推荐写法:有副作用的调用 → io_callback;ordered=True 保证执行顺序与程序顺序一致
from jax.experimental import io_callback

order = []
def write_log(i):
    order.append(int(i))                     # 副作用:写外部状态

def body(i, carry):
    io_callback(write_log, None, i, ordered=True)   # 返回 None 时 result_shape_dtypes 传 None
    return carry + i

total = jax.jit(lambda: jax.lax.fori_loop(0, 5, body, 0))()
jax.effects_barrier()
print("写入顺序:", order, "| total =", total)
assert order == [0, 1, 2, 3, 4] and int(total) == 10
写入顺序: [0, 1, 2, 3, 4] | total = 10

三种回调在 jit / vmap / grad / DCE 下的行为(逐项实测)

下一个单元把每种回调放进各种变换里逐个试一遍,得到下面这张表,并用 assert 把结论固定下来。“输出未被使用时”这一列说的是 DCE(死代码消除):回调的输出没被用到时,jit 会不会把这次调用当成死代码删掉。

回调jitvmapgrad输出未被使用时
jax.debug.callback✅✅(逐样本调用 B 次)✅(只看到 primal 值,对梯度无贡献)保留(有副作用)
io_callback(ordered=False)✅✅(逐样本调用 B 次)❌ ValueError:不支持 JVP保留
io_callback(ordered=True)✅❌ ValueError:不能 vmap 有序回调❌ ValueError保留
pure_callback(不传 vmap_method)✅❌ NotImplementedError❌ ValueError:提示用 custom_jvp被删除,不执行
pure_callback(vmap_method="expand_dims")✅✅(一次调用,收到整个批次)❌ ValueError被删除
# 逐项实测上表
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(默认)":      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调用次数"] = 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["未使用时执行次数"] = len(calls)
    table[name] = row

print(f"{'回调':28s}{'jit':6s}{'vmap':22s}{'vmap调用次数':10s}{'grad':12s}{'未使用时执行次数'}")
for name, r in table.items():
    print(f"{name:28s}{r['jit']:6s}{r['vmap']:22s}{r['vmap调用次数']!s:14s}{r['grad']:12s}{r['未使用时执行次数']}")

assert all(r["jit"] == "OK" for r in table.values())
assert table["debug.callback"]["vmap"] == "OK" and table["debug.callback"]["vmap调用次数"] == 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(默认)"]["vmap"] == "NotImplementedError"
assert table["pure_callback(默认)"]["grad"] == "ValueError"
assert table["pure_callback(expand_dims)"]["vmap"] == "OK" and table["pure_callback(expand_dims)"]["vmap调用次数"] == 1
assert table["debug.callback"]["未使用时执行次数"] == 1 and table["io_callback"]["未使用时执行次数"] == 1
assert table["pure_callback(默认)"]["未使用时执行次数"] == 0      # 纯回调被当作死代码删除
回调                          jit   vmap                  vmap调用次数  grad        未使用时执行次数
debug.callback              OK    OK                    4             OK          1
io_callback                 OK    OK                    4             ValueError  1
io_callback(ordered)        OK    ValueError            0             ValueError  1
pure_callback(默认)           OK    NotImplementedError   0             ValueError  0
pure_callback(expand_dims)  OK    OK                    1             ValueError  0

🔍 补充

🟢 必读

  • 怎么选:只想看一眼中间值,用 jax.debug.print 或 jax.debug.callback。要写文件、发指标,或者依赖执行顺序,用 io_callback(ordered=True)。纯计算并且需要返回值,用 pure_callback。pure_callback 是唯一一种允许 JAX 删除、合并或重复执行的回调,因为 JAX 相信它是纯的。它要是带了副作用,结果就没法预测(见 5.4 的 DCE 验证)。
  • debug.callback 在 grad 下只看得到 primal 值:它不参与微分,看不到切线,也看不到余切。想打印梯度,就对 jax.grad 的结果调用 jax.debug.print。
  • 异步:回调和设备计算是异步的。在 notebook 里,打印内容可能“串”到后面的单元去,用 jax.effects_barrier() 可以等所有回调都执行完。
  • 想让回调可微:三种回调都不支持自动微分,要配合 jax.custom_jvp 或 jax.custom_vjp 使用,见 5.4。

5.4 使用 jax.pure_callback 进行无副作用调用

5.4 使用 jax.pure_callback 进行无副作用调用

如果要把一个纯计算的外部 Python 函数,接进需要 jit、vmap 或 grad 的计算图,就用 jax.pure_callback。

纯粹性约定:被包装的函数必须是函数式纯的。也就是说,它是确定的(同样的输入总是得到同样的输出),而且没有副作用(不改外部状态、不做 I/O,结果只取决于传进来的参数)。JAX 完全相信你遵守了这个约定,不会去检查。违反了的话,结果可能是错的,而且通常不会有明确的报错。

工作原理:调用时要给三样东西:回调函数、描述输出结构的 result_shape_dtypes,以及实际参数。追踪阶段不会执行回调,只在 jaxpr 里放一个占位的 pure_callback 方程,把回调当成黑盒。到运行时,才用具体的值去调用它。

result_shape_dtypes 的写法:单个输出用 jax.ShapeDtypeStruct(shape, dtype)。多个输出用元组,每个元素一个 ShapeDtypeStruct。它支持任意 pytree,只要结构和回调的返回值一致就行。

使用场景:接入 SciPy 这类外部库;用纯 Python 实现复杂逻辑;JAX 缺某个运算时临时顶一下。

课程列出的注意事项:① 纯粹性由你自己负责(这一条最重要);② 加速器和主机之间来回切换、传数据有开销,原生 XLA 运算通常快得多;③ 和变换的兼容性;④ 闭包或依赖外部模块的函数可能不好序列化。

⚠️ 注意:很容易以为 pure_callback 和 jit、vmap、grad 都能直接配合,比如觉得“回调里算的是 sin(x)*y,本身可微,所以能直接 jax.grad”。在 JAX 0.12 里并不是这样:

  1. vmap:vmap_method 默认是 None,这时 vmap 直接抛 NotImplementedError,必须显式指定。0.12 的合法取值有 "sequential"、"sequential_unrolled"、"expand_dims"、"broadcast_all" 和 "legacy_vectorized"。
  2. grad:不管回调里的数学可不可微,都会抛 ValueError: Pure callbacks do not support JVP。JAX 看不到 NumPy 内部在算什么,必须用 jax.custom_jvp 或 jax.custom_vjp 手动提供导数。
  3. 旧参数 vectorized=True 在 0.12 里已经不是 pure_callback 的参数了。它会被当成 **kwargs 原样转发给回调函数,结果运行时报 TypeError: ... unexpected keyword argument 'vectorized'(见 5.10)。
# ✅ 推荐写法:课程主示例(jit 部分),key 用 jax.random.key 生成
def external_pure_python_computation(x, y):
    """纯 Python/NumPy 函数:确定性、无副作用"""
    x = np.asarray(x)                     # 回调收到 jax.Array,转 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")        # 显式声明 vmap 语义(见下文对照)
    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("输出:\n", output)
np.testing.assert_allclose(output, expected, rtol=1e-5)

# vmap:expand_dims 下回调只调用一次,收到 (10, 3, 3) 的整批数据
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)
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)
输出:
 [[ 4.6199546   4.8222537  -0.27386507]
 [ 0.7643099   1.5264554  -1.6242192 ]
 [-0.44588774  2.4433482   2.8974433 ]]
批量输出形状: (10, 3, 3)
# ❌ 旧写法(过时/错误):不指定 vmap_method 直接 vmap;直接对 pure_callback 求 grad
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.
# ✅ 推荐写法:custom_jvp 给 pure_callback 补上导数(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)
    # 导数部分用 jnp 写,它本身可微、可 jit、可 vmap
    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("梯度:\n", gradients)
np.testing.assert_allclose(gradients, np.cos(np.asarray(input_array)) * 2.0 + 1.0, rtol=1e-5)
# 对标量参数 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 也能组合
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)
梯度:
 [[0.8963548 0.1220305 2.8149228]
 [2.9938226 2.9690719 2.1271498]
 [2.7596536 2.7605276 2.5746362]]
# 课程多输出示例:result_shape_dtypes 用元组(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),   # 标量形状为 ()
                        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"多个输出:{res1=}, {res2=}")
assert float(res1) == 20.0 and float(res2) == 6.0
多个输出:res1=Array(20., dtype=float32), res2=Array(6., dtype=float32)

vmap_method 的取值(0.12 实测)

设批大小为 $B$,回调原本处理的输入形状是 $s$:

vmap_method回调被调用几次回调收到的形状什么时候用
"sequential"$B$ 次(lax.scan 逐样本循环)$s$回调只能处理单个样本
"sequential_unrolled"$B$ 次(完全展开的 scan,unroll=B)$s$同上,B 很小时
"expand_dims"1 次被批处理的参数是 $(B, *s)$;未批处理的参数补一个大小为 1 的前导维回调支持 NumPy 广播(最常用)
"broadcast_all"1 次所有参数都广播成 $(B, *s)$回调要求所有输入的批维度一致
"legacy_vectorized"1 次被批处理的参数是 $(B, *s)$,未批处理的参数保持原形状(旧 vectorized=True 的行为)只用于迁移旧代码

不管用哪种方式,批维度都在最前面,回调返回值的第一维也必须是 $B$。

# 逐个试一遍 vmap_method:调用次数与回调收到的形状
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} 调用 {len(seen)} 次,收到 (a, b) 形状 = {seen[0]}  输出 {out.shape}")

assert len(results["sequential"]) == 4 and results["sequential"][0] == ((3,), (3,))
assert results["expand_dims"] == [((4, 3), (1, 3))]       # 未批处理的 b 只多一个 1 维
assert results["broadcast_all"] == [((4, 3), (4, 3))]     # b 被广播到批大小
assert results["legacy_vectorized"] == [((4, 3), (3,))]   # b 保持原形状
try:
    jax.vmap(call_with("bogus"))(jnp.ones((4, 3)))
except ValueError as e:
    print("非法取值 →", str(e)[:120])
sequential           调用 4 次,收到 (a, b) 形状 = ((3,), (3,))  输出 (4, 3)
sequential_unrolled  调用 4 次,收到 (a, b) 形状 = ((3,), (3,))  输出 (4, 3)
expand_dims          调用 1 次,收到 (a, b) 形状 = ((4, 3), (1, 3))  输出 (4, 3)
broadcast_all        调用 1 次,收到 (a, b) 形状 = ((4, 3), (4, 3))  输出 (4, 3)
legacy_vectorized    调用 1 次,收到 (a, b) 形状 = ((4, 3), (3,))  输出 (4, 3)
非法取值 → vmap_method must be on of the allowed methods ['sequential', 'sequential_unrolled', 'expand_dims', 'broadcast_all', 'leg

🔍 补充

🟢 必读

  • DCE 与纯粹性:pure_callback 的输出没被用到时,jit 会把整个调用删掉,回调根本不会执行(见 5.3 的表)。JAX 也可能在变换里重复调用它,比如 sequential 方式的 vmap 会调用 B 次。所以千万不要在 pure_callback 里写日志、计数或做 I/O。
  • 形状和 dtype 必须对得上:返回的形状不对,运行时报 JaxRuntimeError: CpuCallback error。dtype 不对时,有时报错(比如声明 int32 却返回 float),有时又被悄悄转换(关闭 x64 时 float64 会被转成 float32)。稳妥的做法是在回调末尾显式写 .astype(声明的 dtype)。
  • 高阶导数:如果 custom_jvp 里的导数也用 pure_callback 来算,一阶 grad 没问题,但 jax.hessian 会再次碰到“不支持 JVP”。想支持高阶导数,导数函数本身也得可微,比如用 jnp 写,或者再套一层 custom_jvp。下面的单元验证了这一点。
  • 性能:每次调用都要在 XLA 和 Python 之间跨一次,有固定开销,还会挡住融合。下面用计时对比一下。
# 验证:导数本身也用 pure_callback 时,hessian 会失败;把导数也包成 custom_jvp 即可修复
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

# 修复:sin 和 cos 互为导数,两者都定义 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 对角线:", np.diag(H))
hessian(v1) → ValueError: Pure callbacks do not support JVP. Please use `jax.custom_jvp` to use callbacks 
修复后 hessian 对角线: [-0.         -0.32719472 -0.6183698  -0.841471  ]
# 验证:形状不匹配会在运行时报错;输出未被使用时回调不会执行
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("返回形状不符 →", shape_err)
assert shape_err == "JaxRuntimeError"

counter = []
def impure(a):
    counter.append(1)                         # ❌ 在 pure_callback 里做副作用(反例)
    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("输出未被使用时回调执行次数:", len(counter))
assert len(counter) == 0
返回形状不符 → JaxRuntimeError
输出未被使用时回调执行次数: 0
# 计时:pure_callback(NumPy sin)vs 原生 jnp.sin(只打印)
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 | 原生 jnp: {t_nat:.3f} ms | 比值 {t_cb / t_nat:.1f}x")
pure_callback: 2.389 ms | 原生 jnp: 2.824 ms | 比值 0.8x

5.5 JAX 原语简介

5.5 JAX 原语简介

原语是什么:原语(primitive)是 JAX 认得的最基本的原子运算。jnp.add、jnp.dot 这类 NumPy 风格的函数只是 Python 包装,追踪时会被拆成原语,记进 jaxpr。常见的原语有 add、mul、dot_general、reduce_sum、transpose、conv_general_dilated,还有控制流原语 scan、cond。

为什么重要(课程讲了三点):

  1. 变换作用在原语上:jit 把函数追踪成 jaxpr,再交给 XLA 编译;grad 靠的是每个原语的求导规则;vmap 靠的是每个原语的批处理规则。变换从来不看 Python 代码,只看原语。
  2. 编译接口:原语连着高层 API 和编译器。jaxpr 被降低(lower)成编译器 IR,再编译成 CPU/GPU/TPU 的机器码。
  3. 可扩展性:想接入 C++/CUDA 写的专用算法,或者定义一个求导规则很特别的运算,就要定义新原语,并给它配齐各项规则(5.6~5.9)。

原语和函数的区别:jnp.add 是 Python 函数,追踪时它在 jaxpr 里登记的是原语 add。变换和编译需要的所有规则,都挂在原语对象上。

⚠️ 注意:控制流原语的名字是 scan、cond,不是 scan_p、cond_p;jnp.add 登记的原语也叫 add,不叫 add_p。*_p 只是源码里存放原语对象的 Python 变量名,原语自己的名字(.name,也就是 jaxpr 里显示的那个)不带 _p。另外,现在 JAX 的编译流程是 jaxpr → StableHLO(MLIR 方言)→ XLA 编译,并不是 jaxpr 直接转成 XLA HLO,见下面 .lower().as_text() 的输出。

# 查看一个函数分解成哪些原语
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("原语序列:", names)
assert {"sin", "mul", "dot_general", "add", "reduce_sum"} <= set(names)

# 控制流原语的名字
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("控制流原语:", cf_names & {"cond", "scan"})
assert {"cond", "scan"} <= cf_names

import jax.extend.core.primitives as P          # 内置原语对象的半公开入口
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,) }
原语序列: ['sin', 'mul', 'dot_general', 'add', 'reduce_sum']
控制流原语: {'cond', 'scan'}
P.scan_p.name = scan | P.cond_p.name = cond | jax.lax.add_p.name = add
# 原语身上挂着各种规则表:JVP、转置、批处理、lowering
from jax.interpreters import ad, batching, mlir
print("sin 有 JVP 规则     :", P.sin_p in ad.primitive_jvps)
print("mul 有转置规则      :", P.mul_p in ad.primitive_transposes)
# 注意:0.12 的 batching.primitive_batchers 是“只写代理”(只支持 [p] = rule),不能用 in 查询;
# 它把规则包装后存进 batching.fancy_primitive_batchers
try:
    P.sin_p in batching.primitive_batchers
except TypeError as e:
    print("primitive_batchers 不支持 in 查询:", e)
print("sin 有 batching 规则:", 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

# 编译流水线:jaxpr → StableHLO(MLIR 文本)
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 有 JVP 规则     : True
mul 有转置规则      : True
primitive_batchers 不支持 in 查询: argument of type 'PrimitiveBatchersProxy' is not a container or iterable
sin 有 batching 规则: 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>
  }
}

🔍 补充

⚪ 选读:要从零定义自定义原语时

  • 一个原语要完整工作,需要这些规则(后面几节逐个实现):
规则注册位置(0.12)服务的变换缺失时的报错
impl(eager 实现)p.def_impl直接调用(不在 jit 下)Evaluation rule ... not implemented
抽象求值p.def_abstract_eval一切追踪:jit / eval_shape / grad / vmapAbstract evaluation ... not implemented
loweringmlir.register_loweringjit(编译)MLIR translation rule ... not found
JVPad.primitive_jvps[p]jvp / gradDifferentiation rule ... not implemented
转置ad.primitive_transposes[p]grad / vjp(只有线性原语需要)Transpose rule ... not implemented
batchingbatching.primitive_batchers[p]vmapBatching rule ... not implemented
  • batching.primitive_batchers 在 0.12 里是个只写代理:batching.primitive_batchers[p] = rule 会把规则包装成新签名,存进 batching.fancy_primitive_batchers。它不支持 in,也不能读取,要查询就用后者。另外,所有参数都没被批处理时,包装器不会调用你的规则,而是直接 bind。半内部 API 跨版本会变,这就是个典型例子。
  • 这些注册表都是半内部 API。大多数“想接外部代码”的需求,用 pure_callback 或 jax.ffi.ffi_call 再加 custom_vjp 就够了(5.10),用不着自定义原语。

5.6 定义自定义原语

5.6 定义自定义原语

动机:jax.numpy 和 jax.lax 已经覆盖了绝大多数运算。只有下面几种情况才需要新原语:① 性能,要调用预编译、高度优化的内核;② 不支持的运算,现有原语表达不了,或者要用到特殊的硬件功能;③ 对接自定义硬件。课程打了个比方:原语是 JAX 描述计算用的“词汇”,自定义原语就是往词汇表里加一个新词。

定义一个原语要讲清四件事:标识(唯一的名字)、抽象行为(只根据输入的形状和 dtype,推出输出的形状和 dtype)、执行(各后端怎么实现,通常是降低成编译器 IR)、微分(JVP / 转置规则)。

课程的五个步骤:

  1. 创建原语对象,名字会显示在 jaxpr 中。
  2. 写一个给用户调用的包装函数,里面调用 p.bind(*args, **params)。bind 不会马上算出结果,而是交给当前的追踪器(trace)处理:eager 时调用 impl,jit 时记进 jaxpr。会影响输出结构的静态参数(比如 some_parameter)要用关键字参数传给 bind。
  3. 用 def_abstract_eval 注册抽象求值规则(5.7)。
  4. 用 mlir.register_lowering 注册后端 lowering 规则(5.8)。
  5. 注册求导规则(5.9)。

课程还提到:如果原语只是去调用外部代码,很多时候直接在包装函数上用 jax.custom_vjp / jax.custom_jvp 就行,这比在原语上定义规则好维护。自定义原语的代价是:复杂度高、内部 API 跨版本会变、调试困难。所以应该先考虑 jax.lax、回调,或者换个写法。

本节用课程里的 my_custom_op 把完整流程走一遍。它的数学定义取 JVP 注释里提到的 $f(x) = x^2 + c$($c$ = some_parameter):

$$f(x; c) = x^2 + c,\qquad \dot y = 2x\,\dot x$$

⚠️ 注意:在 JAX 0.12 里,from jax.core import Primitive 会报 ImportError,因为 jax.core 已经不再导出 Primitive。推荐的路径是 from jax.extend.core import Primitive。ShapedArray 还在 jax.core.ShapedArray,jax.extend.core 里没有它。

# ❌ 旧写法(过时/错误):from jax.core import Primitive —— 0.12 中 jax.core 不再导出 Primitive
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 有 Primitive:", hasattr(jax.core, "Primitive"), "| 有 ShapedArray:", hasattr(jax.core, "ShapedArray"))
print("jax.extend.core 有 Primitive:", hasattr(jax.extend.core, "Primitive"), "| 有 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 有 Primitive: False | 有 ShapedArray: True
jax.extend.core 有 Primitive: True | 有 ShapedArray: False
# ✅ 推荐写法:步骤 1~3 —— 原语对象、用户函数、impl 与抽象求值
from jax.extend.core import Primitive          # 半内部 API:0.12 下验证过的路径
from jax.core import ShapedArray               # 抽象值类型
from jax.interpreters import mlir, ad, batching

my_custom_op_p = Primitive("my_custom_op")      # 步骤 1:唯一名字

def my_custom_op(x, *, some_parameter):         # 步骤 2:面向用户的函数
    return my_custom_op_p.bind(x, some_parameter=some_parameter)

@my_custom_op_p.def_impl                         # eager 实现:收到具体数组
def _my_custom_op_impl(x, *, some_parameter):
    return x * x + some_parameter

@my_custom_op_p.def_abstract_eval                # 步骤 3:只看形状/dtype,不算数值
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])

# 还没有 lowering → jit 失败;没有 JVP → grad 失败;没有 batching → vmap 失败
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, "→ 意外成功")
    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
# ✅ 推荐写法:步骤 4 —— lowering(详见 5.8;这里用 mlir.lower_fun 把 jnp 实现降低成 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.]
# ❌ 旧写法(过时/错误):JVP 规则里用 ad.Zero.from_value(primal_out) 构造零切线 —— 0.12 中该方法不存在
print("ad.Zero 有 from_value:", hasattr(ad.Zero, "from_value"))
try:
    ad.Zero.from_value(jnp.ones(3))
except AttributeError as e:
    print("AttributeError:", e)
ad.Zero 有 from_value: False
AttributeError: type object 'Zero' has no attribute 'from_value'
# ✅ 推荐写法:步骤 5 —— JVP 规则(ẏ = 2x·ẋ),零切线用 ad.Zero(切线 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:                                  # 输入被当作常数
        tangent_out = ad.Zero(jax.typeof(primal_out).to_tangent_aval())
    else:
        tangent_out = 2 * x * x_dot                             # 用现有原语表达,对 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:batching.defvectorized 一行搞定(batch 维原样透传)
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) 通过")
grad: [2. 4. 6.] | jit(grad): [2. 4. 6.]
vmap / vmap(grad) 通过

⚠️ 注意:ad.primitive_transposes 注册的是转置规则,不是 VJP。JAX 的反向模式 = JVP 线性化 + 转置。只有当 JVP 的切线部分用到了本原语自己时(这时本原语对切线是线性的),才需要转置规则。上面的 JVP 用 2 * x * x_dot(内置的 mul)表达切线,mul 自带转置规则,所以只注册 JVP,jax.grad 就能工作,上面的 assert 已经验证过。需要自己写转置规则的情况见 5.9。


5.7 实现抽象求值规则

5.7 实现抽象求值规则

抽象求值是什么:在 jit、自动微分和 vmap 真正开始计算之前,JAX 要先知道每个运算输出的形状和 dtype,但不算具体数值。它用抽象值(主要是 jax.core.ShapedArray,只有 shape 和 dtype)来追踪程序。

目的:抽象求值相当于原语在“类型和形状层面的函数签名”。追踪时遇到一个原语,JAX 就用(抽象输入, 参数)去调用它,得到抽象输出。课程给了三个理由:

  1. 静态分析:还没编译、没执行,就能发现形状不匹配和类型错误。
  2. 变换兼容:vmap 要靠它推出批处理后的形状,grad 要靠它搭出反向传播的结构。
  3. 编译:XLA 要知道每个运算准确的形状和类型,才能生成代码。

参数(params):如果输出形状取决于非数组参数(卷积的步长、归约的轴等),这些参数就作为 bind 的关键字参数传进来,抽象求值规则也从关键字参数里拿到它们。

在流程里的位置:抽象求值是新原语第一条要写的规则,lowering 和 JVP 都要用到它给出的形状信息。jaxpr 就是靠这些抽象值串起来的。

课程给了两个例子:custom_add_one(加 1,并检查输入个数、类型和浮点 dtype),以及 custom_reduce_sum(按 axis 参数去掉维度)。下面把两个例子都完整跑一遍。

# ✅ 课程示例 1:custom_add_one 的抽象求值(含输入校验)
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 需要恰好一个输入操作数。")
    input_aval = avals[0]
    if not isinstance(input_aval, ShapedArray):
        raise TypeError(f"期望 ShapedArray,得到 {type(input_aval)}")
    if not jnp.issubdtype(input_aval.dtype, jnp.floating):
        raise TypeError(f"custom_add_one 只支持浮点输入,得到 {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 只做抽象求值:得到形状,但 impl 一次都不执行
out_sds = jax.eval_shape(custom_add_one_p.bind, jax.ShapeDtypeStruct((4, 5), jnp.float32))
print("eval_shape:", out_sds, "| impl 调用次数:", len(impl_calls))
assert out_sds.shape == (4, 5) and out_sds.dtype == jnp.float32 and impl_calls == []

# 类型错误在追踪期就被发现(还没编译、没执行)
try:
    jax.eval_shape(custom_add_one_p.bind, jax.ShapeDtypeStruct((3,), jnp.int32))
    caught = False
except TypeError as e:
    caught = True
    print("追踪期 TypeError:", e)
assert caught
eval_shape: ShapeDtypeStruct(shape=(4, 5), dtype=float32) | impl 调用次数: 0
追踪期 TypeError: custom_add_one 只支持浮点输入,得到 int32
# ✅ 课程示例 2:custom_reduce_sum —— 输出形状由参数 axis 决定
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 必须可哈希(tuple)

@custom_reduce_sum_p.def_abstract_eval
def custom_reduce_sum_abstract_eval(aval, *, axis):
    if not isinstance(aval, ShapedArray):
        raise TypeError("输入必须是 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,) }

⚠️ 注意:多输出的原语,光让抽象求值返回 ShapedArray 的元组还不够,还必须先设置 p.multiple_results = True。不然抽象求值返回元组时会报 ValueError: ... should return a tuple or a list iff ... multiple_results。同时,impl 和 lowering 也要按多输出来写(lower_fun(..., multiple_results=True))。

# 多输出原语:minmax(x) -> (min, max)
minmax_bad = Primitive("minmax_bad")                    # 忘了设 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

🔍 补充

⚪ 选读:要给自定义原语写抽象求值规则时

  • 抽象求值里只能有形状逻辑:它收到的是 ShapedArray,不是数组。在里面做 float(x)、np.asarray(x) 这类取值操作都会失败。凡是要知道“值”才能决定的东西(比如输出形状),只能从静态参数里来。
  • 参数必须可哈希:bind 的关键字参数会成为 jaxpr 方程的一部分,还参与 jit 缓存,所以要用 tuple,不能用 list 或数组。上面的 custom_reduce_sum 先把 axis 整理成排好序的 tuple,意思相同的写法就能命中同一个缓存。
  • 调试抽象求值,jax.eval_shape 很好用:它只追踪不执行,可以很快检查形状推导和报错。

5.8 为后端(CPU/GPU/TPU)实现转换(lowering)规则

5.8 为后端实现转换规则

为什么需要:jit 会把函数交给 XLA 编译。遇到原语时,JAX 得把它翻译成编译器看得懂的低层指令。没有这条“转换(lowering)规则”,XLA 就不知道该给自定义运算生成什么代码。不同后端上最好的实现可能不一样,所以可以按平台分别注册。

课程描述的机制:转换规则接收 XLA 值,用 XLA builder 构建 HLO 图(builder.ConstantF32Scalar、builder.Add、builder.Mul、builder.DotGeneral、builder.Reshape…),再用 xla_client.register_translation(prim, rule, platform=...) 注册。

各后端的考虑:CPU 上 XLA 可以利用多核和向量指令(AVX 等)。GPU 上如果没有直接对应的 HLO,可以通过 CustomCall 嵌入 CUDA 内核。TPU 上要考虑矩阵单元(MXU),通常把计算组织成大矩阵乘法。

课程示例:ReLU $\mathrm{relu}(x) = \max(0, x)$。先从操作数取出形状和 dtype,构造一个同形状的零张量,再做逐元素 Max。XLA 知道怎么高效编译 Max,所以一条规则就够所有后端用。

更复杂的情况:组合多个现有的 HLO 运算;给 GPU 写专用内核,再通过 CustomCall 调用;搞清楚 XLA 融合对生成代码有什么影响。调试时一般要看生成的 IR。

⚠️ 注意:from jax.lib import xla_client、xla_client.register_translation 和 XLA builder(builder.Add、builder.GetShape 等)都属于旧的 XLA 翻译规则 API。在 JAX 0.12 里,jax.lib.xla_client 导入会失败,jax.interpreters.xla 里也没有 register_translation(只剩 apply_primitive 等少数几个成员)。

现在的写法是 MLIR lowering:用 jax.interpreters.mlir.register_lowering(prim, rule, platform=None|'cpu'|'cuda'|'rocm'|'tpu') 注册。规则的签名是 rule(ctx, *mlir_values, **params) -> list[ir.Value],生成的是 StableHLO。最简单的写法是 mlir.lower_fun(jnp实现, multiple_results=...)。想精细控制,就用 jax.extend.mlir.dialects.stablehlo 手写 op。要调用外部 C++,就用 jax.ffi.ffi_lowering(5.10,练习 4)。

# ❌ 旧写法(过时/错误):xla_client.register_translation + XLA builder —— 0.12 中已不存在
try:
    from jax.lib import xla_client
except ImportError as e:
    print("ImportError:", short_path(e))
from jax.interpreters import xla
print("jax.interpreters.xla 公开成员:", [a for a in dir(xla) if not a.startswith("_")])
print("有 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)
jax.interpreters.xla 公开成员: ['Backend', 'apply_primitive', 'canonicalize_dtype_handlers']
有 register_translation: False
# ✅ 推荐写法 A:mlir.lower_fun —— 用 jnp 写出语义,由 JAX 自动降低成 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))

# 通用规则(所有平台):用 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>
  }
}

# ✅ 推荐写法 B:手写 StableHLO(对应课程的 ReLU 示例),并只为 'cpu' 平台注册
from jax.extend.mlir.dialects import stablehlo as hlo

def my_relu_lowering_cpu(ctx, x_operand):
    aval_out, = ctx.avals_out                                       # 抽象求值给出的输出 aval
    out_type = mlir.aval_to_ir_type(ctx.module_context, aval_out)   # 0.12:需要传 module_context
    zero = mlir.ir_constant(np.zeros((), aval_out.dtype))           # 标量 0
    zeros = hlo.broadcast_in_dim(out_type, zero, mlir.dense_int_array([]))   # 广播成同形状
    return [hlo.maximum(zeros, x_operand)]                           # 必须返回 Value 列表

mlir.register_lowering(my_relu_p, my_relu_lowering_cpu, platform="cpu")   # 平台专用规则优先于通用规则

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>
  }
}

🔍 补充

⚪ 选读:要手写 lowering、或按平台注册不同实现时

  • platform 的取值:None 表示通用规则,平台专用规则('cpu'、'cuda'、'rocm'、'tpu')的优先级更高。GPU 上的平台名是 'cuda' / 'rocm',不是 'gpu'。
  • lowering 不等于 impl:jit 走 lowering,eager 走 def_impl。两边实现不一致的话,同一个函数加不加 jit 结果会不一样,这种 bug 很难发现。省事的做法是让 impl 直接复用 lowering:p.def_impl(partial(xla.apply_primitive, p)),eager 调用时就临时 jit 这一个原语(练习 4 会用到)。
  • ctx 里有用的字段:ctx.avals_in / ctx.avals_out(抽象值)、ctx.module_context(构造类型时要用)、ctx.platforms。

5.9 制定自定义原语的求导规则

5.9 制定自定义原语的求导规则

为什么需要:只有抽象求值和 lowering 的原语,还不能参与自动微分。对它用 jax.grad、jax.jvp 或 jax.vjp 会报错,因为 JAX 不知道导数该怎么传。

前向模式(JVP):设 $y = f(x)$,雅可比矩阵 $J = \partial f / \partial x$,JVP 计算

$$\dot y = J\,\dot x .$$

规则接收 (primals, tangents),返回 (primal_out, tangent_out)。切线的形状和 dtype 跟对应的 primal 一样。被当作常数的输入,它的切线是符号零 ad.Zero。

反向模式(VJP):计算

$$\bar x^{\top} = \bar y^{\top} J .$$

jax.grad 走的就是反向模式。课程讲的 VJP 是“前向保存残差,backward_pass 计算输入余切”的结构,也就是 custom_vjp 的写法。

课程示例 custom_scale:$y = x \cdot c$($c$ 是标量 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$ 是标量,所以 $\bar c$ 要对所有元素求和。用 $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)$$

JVP 的预期值:$\dot x = [0.1, 0.2, 0.3]$、$\dot c = 0.5$ 时,$\dot y = [0.7, 1.4, 2.1]$。

验证方法(课程):① 数值检查,用 jax.test_util.check_grads 拿有限差分和自定义规则做比较;② 检查 JVP 和 VJP 是否一致。

⚠️ 注意:旧资料里常见 custom_scale_p.def_jvp / def_vjp 注册规则、在抽象求值里调用 jax.core.get_aval、用 ad.Zero.zero_if_zero 构造零的写法。在 JAX 0.12 里,Primitive 根本没有 def_jvp / def_vjp 方法,jax.core.get_aval 和 ad.Zero.zero_if_zero 也都不存在。原语层面正确的做法是:

  • JVP:ad.primitive_jvps[p] = jvp_rule;
  • 反向模式:不注册 VJP。JAX 先用 JVP 线性化,再转置切线计算。如果 JVP 的切线部分用到了本原语,就要注册 ad.primitive_transposes[p] = transpose_rule;
  • “前向 + 残差 + backward_pass”这种结构属于 jax.custom_vjp,适合用在包装函数这一层(见 5.10)。

还有一点:写成 custom_scale_p.bind(x, factor=factor) 的话,factor 就成了参数;而抽象求值的签名 (x_abs, factor_abs) 又把它当成第二个操作数,两边对不上。正确做法是:factor 要求导,就必须作为操作数(位置参数)传给 bind。参数(params)是静态的,没法对它求导。

# ❌ 旧写法(过时/错误):Primitive.def_jvp / def_vjp、jax.core.get_aval、ad.Zero.zero_if_zero 均不存在;bind 参数与抽象求值签名矛盾
print("Primitive 有 def_jvp:", hasattr(Primitive, "def_jvp"), "| 有 def_vjp:", hasattr(Primitive, "def_vjp"))
print("jax.core 有 get_aval:", hasattr(jax.core, "get_aval"), "| ad.Zero 有 zero_if_zero:", hasattr(ad.Zero, "zero_if_zero"))
print("jax.ad 有 Zero:", hasattr(jax.ad, "Zero"), "(jax.ad 现在是另一个新命名空间)")

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 作为参数传入
except TypeError as e:
    print("bind(x, factor=...) → TypeError:", str(e).splitlines()[0][:120])
Primitive 有 def_jvp: False | 有 def_vjp: False
jax.core 有 get_aval: False | ad.Zero 有 zero_if_zero: False
jax.ad 有 Zero: False (jax.ad 现在是另一个新命名空间)
def_jvp → AttributeError: 'Primitive' object has no attribute 'def_jvp'
bind(x, factor=...) → TypeError: <lambda>() got an unexpected keyword argument 'factor'
# ✅ 推荐写法:custom_scale 原语 —— x 和 factor 都是操作数
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 必须是标量,得到形状 {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·ċ —— 第一项用了本原语本身(对 ẋ 线性),因此反向模式需要转置规则
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)

# 对 factor 求导:切线路径只经过内置 mul,无需转置规则 → 已经可以运行
print("dL/dfactor(尚无转置规则):", jax.grad(loss_fn, argnums=1)(x_val, factor_val))
assert float(jax.grad(loss_fn, argnums=1)(x_val, factor_val)) == 56.0
# 对 x 求导:需要转置 custom_scale(ẋ, c)
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(尚无转置规则): 56.0
dL/dx → Transpose rule (for reverse-mode differentiation) for 'custom_scale' not implemented
# ✅ 转置规则:对线性参数 x(UndefinedPrimal)给出余切 x̄ = ȳ·c;非线性参数 factor 返回 None
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 仍用本原语表达
    return ct_x, None                            # 与操作数一一对应

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

# 课程的验证方法:check_grads(一阶 + 二阶,前向 + 反向)
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 通过!")
dL/dx = [ 8. 16. 24.] | dL/dfactor = 56.0
check_grads 通过!

🔍 补充:batching 规则

⚪ 选读:自定义原语需要支持 vmap 时

想让原语支持 vmap,还要注册 batching.primitive_batchers[p] = rule。规则的签名是 rule(batched_args, batch_dims, **params) -> (out, out_dim)。其中 batch_dims[i] 是第 i 个参数的批维度位置,None 表示这个参数没有被批处理。规则要做的事,就是用“不带 vmap 的计算”一次处理完整批数据,再告诉 JAX 输出的批维度在哪。

custom_scale 分两种情况:

  • factor 没有被批处理:$y = x\cdot c$ 本身是逐元素的,直接 custom_scale(x, c),输出批维度与 x 相同。
  • factor 被批处理(每个样本的 $c_b$ 不同):把批维度挪到最前面,用广播乘法算 $y_b = x_b \cdot c_b$,输出的批维度是 0。
def _custom_scale_batch(batched_args, batch_dims):
    x, factor = batched_args
    xd, fd = batch_dims
    if fd is None:                                   # factor 未批处理
        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)))
# vmap ∘ grad ∘ jit 组合:每个样本 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 通过 jit / jvp / grad / vmap / vmap(grad) / check_grads")
print(jax.make_jaxpr(jax.grad(loss_fn))(x_val, factor_val))
custom_scale 通过 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,) }

🔍 补充

⚪ 选读:要给自定义原语写求导规则时

  • 为什么只要 JVP + 转置:JAX 的反向模式是 vjp = transpose(linearize(jvp))。每个原语只要写好“前向导数”和“线性部分的转置”,前向、反向、高阶导数(比如 check_grads(order=2))就都能自动拼出来。这是原语层面的设计,比给每个原语分别写 VJP 更容易组合。
  • 转置规则的参数:线性的输入以 ad.UndefinedPrimal 的形式出现,用 ad.is_undefined_primal 判断,它的 .aval 给出形状。非线性的输入就是普通值。返回值和操作数一一对应,非线性输入的位置返回 None。
  • 零切线:ad.Zero(aval) 是符号零,不占内存。规则里要先判断 type(t) is ad.Zero,别直接拿它做算术。需要真正的零数组时,用 ad.instantiate_zeros(t)。
  • 什么时候不用原语:如果只是想给外部代码(回调或 FFI)配上梯度,用 custom_jvp / custom_vjp 包住用户函数就够了(5.4、5.10),根本不用碰 ad.primitive_*。

5.10 实践:整合 C++ 函数

5.10 实践:整合 C++ 函数

JAX/XLA 已经高度优化了,但有时还是需要复用现成的 C++ 库,或者自己实现对性能要求很高的底层运算。课程介绍了两种方法:简单的回调集成,和完整的自定义原语。动手示例用的是前者(ctypes + pure_callback),后者只做了概述。

场景:一个逐元素函数(假装它是一段复杂的 C++ 遗留代码)

$$f(x) = x^2 + 10,\qquad f'(x) = 2x$$

课程步骤:

  1. C++ 实现:用 extern "C" 防止 C++ 的名称修饰(name mangling),这样 ctypes 才能按名字找到符号;对 double 数组逐元素计算 input[i]*input[i] + 10.0。
  2. 编译为共享库:Linux 用 g++ -shared -fPIC -o custom_op.so custom_op.cpp,macOS 用 g++ -shared -o custom_op.dylib custom_op.cpp,Windows 用 MinGW 或 MSVC 的对应命令。
  3. ctypes 封装:用 ctypes.CDLL 加载库,设置 argtypes/restype。NumPy 包装函数负责转成 float64、保证内存连续、分配输出,再把指针传给 C 函数。预期结果是 [1,2,3] → [11,14,19]。
  4. pure_callback 集成:这个 C++ 函数在数学上是纯的,适合用 pure_callback。课程在 eager 和 jit 下测试(预期 [1,2,3,4] → [11,14,19,26]),并指出 jax.grad 会失败,因为自动微分看不透 C++ 代码。
  5. 另一种方案(只做概述):定义 Primitive → 抽象求值 → 各后端 lowering(提到 xla_client.ops.CustomCall 和 xla.register_translation)→ JVP/转置 → 用户函数 bind。
  6. 总结:回调适合纯函数、并且不需要穿过外部代码求梯度的场景。自定义原语集成得最紧密,也支持自动微分,但复杂度高。

本节先按课程的思路,用 0.12 的写法把 ctypes 方案跑通。然后用现在推荐的 XLA FFI(jax.ffi),接入一个真正编译出来的 C++ RMSNorm 内核,并补上梯度和 vmap。

⚠️ 注意:这段集成代码如果照旧写法来写,在 0.12 上有 4 个地方会出问题:

  1. from jax.experimental import host_callback as hcb 直接报 ImportError。哪怕导入之后根本没用到它,整段脚本也会在 import 阶段就失败;
  2. from jax.experimental import pure_callback 报 ImportError,正确的入口是 jax.pure_callback;
  3. pure_callback(..., vectorized=True):vectorized 已经不是 pure_callback 的参数了,它会被当成 kwargs 转发给回调,运行时报错。应该改成 vmap_method="expand_dims" 或 "broadcast_all";
  4. 默认没开 x64 时,jnp.arange(1.0, 5.0, dtype=jnp.float64) 会被截断成 float32(并给出 UserWarning),不能当成 float64 用。

xla_client.ops.CustomCall、xla.register_translation(以及更早流行的、用 pybind11 导出 PyCapsule 的写法)是旧的 custom call 接入方式。0.12 推荐用 XLA FFI:C++ 端用 xla/ffi/api/ffi.h 里的 XLA_FFI_DEFINE_HANDLER_SYMBOL 定义类型安全的 handler;Python 端用 jax.ffi.pycapsule + jax.ffi.register_ffi_target 注册,再用 jax.ffi.ffi_call 调用,不需要 pybind11。

%%writefile cpp/custom_op.cpp
// 步骤 1:逐元素 f(x) = x^2 + 10;extern "C" 防止名称修饰,便于 ctypes 调用
#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;
    }
}
}
Overwriting cpp/custom_op.cpp
# 课程步骤 2:编译为共享库(本机 clang++;在单元里用 subprocess 执行,失败时如实打印错误)
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"编译失败:{src}")
    print("编译成功 →", out.name, f"({out.stat().st_size} 字节)")
    return out

CXX = shutil.which("clang++") or shutil.which("g++") or "c++"   # macOS 用 clang++;Linux / Colab 上通常是 g++
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
编译成功 → libcustom_op.so (16816 字节)
# 课程步骤 3:ctypes 封装(库路径用绝对路径)
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 + 连续内存
    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 封装测试:", 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 封装测试:
 [11. 14. 19.]
# ❌ 旧写法(过时/错误):host_callback / jax.experimental.pure_callback 导入失败;float64 被截断;vectorized=True 被转发给回调
for stmt in ["from jax.experimental import host_callback as hcb",
             "from jax.experimental import pure_callback"]:
    try:
        exec(stmt); print(stmt, "→ 成功")
    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("请求 float64,实际 dtype =", x_jax.dtype, "| 警告:", str(w[0].message)[:70] if w else None)

import logging
_cb_logger = logging.getLogger("jax._src.callback")   # 回调失败时 JAX 会把整段 traceback 打到日志里,这里先关掉
_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__, "| 根因: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
请求 float64,实际 dtype = float32 | 警告: Explicitly requested dtype float64 requested in arange is not availabl
vectorized=True → JaxRuntimeError | 根因:custom_op_numpy() got an unexpected keyword argument 'vectorized'
# ✅ 推荐写法: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),          # 回到声明的 dtype
        jax.ShapeDtypeStruct(x.shape, x.dtype), x,
        vmap_method="expand_dims")                            # 逐元素运算:整批一次调用即可

@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.]

现代做法:XLA FFI(jax.ffi)接入 C++ RMSNorm

RMSNorm 对最后一维做归一化($N$ 是最后一维的长度):

$$s = \Big(\tfrac{1}{N}\sum_{i} x_i^2 + \varepsilon\Big)^{-1/2},\qquad y_i = s\,x_i$$

反向传播时($\bar y$ 是输出余切),由 $\partial s/\partial x_j = -s^3 x_j / N$ 可以推出:

$$\bar x_j = s\,\bar y_j - \frac{s^3}{N}\, x_j \sum_i x_i\,\bar y_i$$

FFI 的流程:

  1. C++ 端:写一个普通函数 ffi::Error Impl(属性..., ffi::Buffer<T> 输入..., ffi::ResultBuffer<T> 输出...),再用 XLA_FFI_DEFINE_HANDLER_SYMBOL(符号名, Impl, ffi::Ffi::Bind().Attr<float>("eps").Arg<...>().Ret<...>()) 声明签名。调用时 XLA 会检查类型,并把形状信息(dimensions())一起传给 C++。
  2. 编译:头文件在 jax.ffi.include_dir() 下面。只要头文件,不用链接任何库。
  3. Python 端:ctypes.cdll.LoadLibrary 加载 → jax.ffi.pycapsule(lib.符号) 包装函数指针 → jax.ffi.register_ffi_target(名字, 胶囊, platform="cpu") 注册。
  4. 调用:jax.ffi.ffi_call(名字, 输出 ShapeDtypeStruct, vmap_method=...)(*输入, **属性)。它生成的是 StableHLO 的 custom_call,天然支持 jit,也不会像回调那样在 Python 里走一个来回。
  5. 梯度:ffi_call 不可微,要用 jax.custom_vjp,配一个前向内核(多输出一个残差 $s$)和一个反向内核。
%%writefile cpp/rms_norm.cc
// XLA FFI 版 RMSNorm:前向、带残差的前向、反向,共三个 handler
#include <cmath>
#include <cstdint>
#include <utility>
#include "xla/ffi/api/ffi.h"

namespace ffi = xla::ffi;

// 对一行长度为 n 的数据做 RMSNorm,返回缩放因子 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;
}

// 反向: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];
}

// 返回 (元素总数, 最后一维长度);把前面所有维度当作“行”
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();
}

// 声明 handler 签名:属性 eps(float32)、输入缓冲区、输出缓冲区
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>>());
Overwriting cpp/rms_norm.cc
# 编译:头文件路径来自 jax.ffi.include_dir(),无需链接任何库
print("FFI 头文件目录:", 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()])

# 加载并注册三个 FFI target
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("已注册: rms_norm, rms_norm_fwd, rms_norm_bwd")
FFI 头文件目录: <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
编译成功 → librms_norm.so (93664 字节)
已注册: rms_norm, rms_norm_fwd, rms_norm_bwd
# ffi_call:jit 下运行,并与 jnp 实现、NumPy float64 实现对拍
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))   # 属性 dtype 必须与 C++ 的 float 一致

def rms_norm_ref(x, eps=1e-5):                  # 纯 jnp 参考实现
    return x / jnp.sqrt(jnp.mean(x ** 2, axis=-1, keepdims=True) + eps)

def rms_norm_np64(x, eps=1e-5):                 # NumPy float64 参考
    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("与 jnp 最大误差:", float(jnp.max(jnp.abs(y_ffi - rms_norm_ref(x)))),
      "| 与 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)          # eager 也能直接调用

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
与 jnp 最大误差: 4.76837158203125e-07 | 与 NumPy f64 最大误差: 3.514383175584612e-07
%0 = stablehlo.custom_call @rms_norm(%arg0) {backend_config = "", mhlo.backend_config = {eps = 9.99999974E-6 : f32}, ope
# 验证两个易错点:属性 dtype 不匹配;ffi_call 默认既不能 vmap 也不能求导
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_method 的 vmap →", 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_method 的 vmap → 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:前向内核额外输出残差 s,反向内核计算 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)                                   # 残差:每行的 s 和输入 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))       # 权重作为参数传入(避开闭包常量 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("梯度最大误差:", float(jnp.max(jnp.abs(g_ffi - g_ref))))
np.testing.assert_allclose(g_ffi, g_ref, rtol=1e-4, atol=1e-5)

# 有限差分检查(反向模式,一阶)
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) 通过")
梯度最大误差: 4.76837158203125e-07
check_grads(rev) 通过
# ✅ vmap:ffi_call 用 vmap_method="broadcast_all"(内核本来就把前导维当“行”处理)
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)
# 批维度不在最前面: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:每个样本都走 fwd/bwd 两个 C++ 内核
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 通过 jit / grad / vmap / vmap(grad),形状:", vg.shape)
FFI RMSNorm 通过 jit / grad / vmap / vmap(grad),形状: (5, 8, 64)
# 计时:FFI 内核 vs pure_callback(NumPy) vs 纯 jnp(只打印比较结果)
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++ 标量循环): {t_ffi:.3f} ms | pure_callback(NumPy): {t_cb:.3f} ms | 纯 jnp(XLA): {t_jnp:.3f} ms")
FFI(C++ 标量循环): 1.040 ms | pure_callback(NumPy): 0.972 ms | 纯 jnp(XLA): 0.079 ms

🔍 补充

⚪ 选读:要把 C++ / CUDA 代码接进 JAX 时

  • FFI 和 pure_callback 怎么选:pure_callback 每次都要回到 Python(受 GIL 限制,还会挡住融合),而且只能在主机上跑。FFI 生成的是 custom_call,XLA 直接调用 C++ 函数指针,不经过 Python,也能写 GPU 内核(在 handler 里拿到 CUDA stream)。性能敏感的场景用 FFI,原型阶段用 pure_callback。不过要看清开销到底花在哪。在本机测下来:对 (8, 64) 这样的小数组,FFI 比 pure_callback 快一个数量级,因为每次进出 Python 的固定开销占了大头;对上面 (4096, 256) 的大数组,两者耗时差不多,因为 NumPy 本身就是向量化的 C 代码,固定开销被摊薄了。
  • C++ 写得朴素,不一定比 XLA 快:上面的计时里,用标量循环写的 C++ 不一定赢得了 XLA 融合后的向量化代码。FFI 真正的价值在于复用现成的高性能库,或者实现 XLA 表达不了的算法。
  • 属性类型要严格对上:Attr<float> 要求传 np.float32(...)。Python 的 float 会被当成 F64,而且要到执行时才报 Wrong scalar data type。整数属性也一样,要用 np.int32 或 np.int64。
  • vmap_method 怎么选:内核能处理任意前导维,就用 "broadcast_all";只能处理单个样本,就用 "sequential"。不指定的话,vmap 直接报错。
  • 符号可见性:XLA_FFI_DEFINE_HANDLER_SYMBOL 生成的是 extern "C" 符号,所以 ctypes 能按名字找到它,不用再自己写 extern "C"。
  • 想把 FFI 做成“原语”:mlir.register_lowering(p, jax.ffi.ffi_lowering("rms_norm"), platform="cpu") 可以直接拿 FFI target 当自定义原语的 lowering,见练习 4。

动手练习

下面有 4 道练习。每道题都配了「✍️ 作答区」和「🧪 自检」,参考答案放在章末附录,里面用 assert 验证。建议先自己写,再去对照参考答案。

练习 1(原语的线性规则):给 5.7 的 custom_reduce_sum_p 补上 JVP、转置和 batching 规则,让 jax.grad、jax.vmap 和 jax.jit 都能通过,并与 jnp.sum 对拍。提示:求和是线性运算,JVP 是 $\dot y = \mathrm{sum}(\dot x)$;转置就是把余切广播回输入形状;batching 时注意,axis 要跟着批维度的位置平移。

自检接口:为 5.7 的 custom_reduce_sum_p 注册 ad.primitive_jvps、ad.primitive_transposes、batching.primitive_batchers 三条规则,之后用 custom_reduce_sum(x, axis) 调用。

练习 2(SciPy + pure_callback + custom_jvp):用 scipy.special.gammaln 和 pure_callback 实现 $\log\Gamma(x)$,导数 $\psi(x)$ 用 scipy.special.digamma(也用回调)。要求支持 vmap(选对 vmap_method)和一阶 grad,并与 jax.scipy.special.gammaln 及其梯度对拍。

自检接口:gammaln_cb(x):用 jax.custom_jvp 定义、内部经 jax.pure_callback 调 SciPy 的 $\log\Gamma$,逐元素作用于任意形状的 float32 数组。

练习 3(C++ FFI 逐元素内核 + 梯度):用 XLA FFI 写一个 C++ SiLU 内核 $\mathrm{silu}(x) = x\,\sigma(x)$,同时写它的反向内核 $\bar x = \bar y\,\sigma(x)\,(1 + x(1-\sigma(x)))$。编译、注册,用 custom_vjp 接上梯度,在 jit、grad 和 vmap 下与 jax.nn.silu 对拍。

自检接口:silu_ffi(x):用 jax.custom_vjp 定义,前向和反向各调用一个经 jax.ffi 注册的 C++ 内核,作用于 float32 数组。

练习 4(FFI 作为原语的 lowering):新建原语 rms_norm_prim,用 jax.ffi.ffi_lowering("rms_norm") 作为它的 CPU lowering,eps 作为参数(params)传入。再补上 eager impl(复用 lowering)和 batching 规则,在 eager、jit 和 vmap 下与 rms_norm_np64 对拍。

自检接口:原语 rms_norm_prim_p,用 rms_norm_prim_p.bind(x, eps=np.float32(eps)) 调用,对最后一维做 RMSNorm;支持 eager、jit、vmap(任意 in_axes)。

# ✍️ 作答区:练习 1
# 在这里写你的代码,然后运行下面的「🧪 自检」。参考答案在本章末尾的「附录:练习参考答案」。
# 🧪 自检:练习 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 规则", custom_reduce_sum_p in ad.primitive_jvps),
                            ("转置规则", custom_reduce_sum_p in ad.primitive_transposes),
                            ("batching 规则", _has_batching_rule())] if not ok]
if _missing:
    print(f"⏳ 练习 1 尚未作答:custom_reduce_sum_p 还没注册 {'、'.join(_missing)}")
else:
    _a = jax.random.normal(jax.random.key(21), (3, 4, 5))
    _ct = jax.random.normal(jax.random.key(22), (4,))
    # 线性:前向、JVP
    np.testing.assert_allclose(custom_reduce_sum(_a, (0, 2)), np.asarray(_a).sum((0, 2)), rtol=1e-5, err_msg="前向求和不对")
    _, _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 不对:切线应为 sum(ẋ)")
    # 转置:把余切广播回输入形状
    _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="转置规则不对:应把余切广播回输入形状")
    _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) 结果不对")
    # batching:批维度在不同位置
    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}) 结果不对:axis 要按批维度位置平移")
    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) 结果不对")
    check_grads(lambda v: custom_reduce_sum(v, (1,)), (_a,), order=2, modes=["fwd", "rev"], eps=1e-2)
    print("✅ 练习 1 自检通过")
⏳ 练习 1 尚未作答:custom_reduce_sum_p 还没注册 JVP 规则、转置规则、batching 规则
# ✍️ 作答区:练习 2
# 在这里写你的代码,然后运行下面的「🧪 自检」。参考答案在本章末尾的「附录:练习参考答案」。
# 🧪 自检:练习 2
try:
    gammaln_cb
except NameError as e:
    print(f"⏳ 练习 2 尚未作答:{e}")
else:
    import scipy.special as _sps
    assert isinstance(gammaln_cb, jax.custom_jvp), "导数要用 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)), "前向应通过 jax.pure_callback 调用 SciPy"
    np.testing.assert_allclose(jax.jit(gammaln_cb)(_x), _sps.gammaln(_x64), rtol=1e-5, atol=1e-6, err_msg="jit 下前向值不对")
    _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="梯度应为 digamma(x)")
    np.testing.assert_allclose(jax.vmap(_gfn)(_x), _sps.digamma(_x64), rtol=1e-4, atol=1e-5, err_msg="vmap(grad) 结果不对:检查 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="嵌套 vmap 下前向值不对")
    print("✅ 练习 2 自检通过")
⏳ 练习 2 尚未作答:name 'gammaln_cb' is not defined
# ✍️ 作答区:练习 3
# 在这里写你的代码,然后运行下面的「🧪 自检」。参考答案在本章末尾的「附录:练习参考答案」。
# 🧪 自检:练习 3
try:
    silu_ffi
except NameError as e:
    print(f"⏳ 练习 3 尚未作答:{e}")
else:
    assert isinstance(silu_ffi, jax.custom_vjp), "梯度要用 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, "前向应调用 FFI 内核(lowering 里没有 custom_call)"
    assert _bwd_hlo.count("stablehlo.custom_call") >= 2, "反向也应调用 C++ 内核:value_and_grad 的 lowering 里应有前向、反向两个 custom_call"
    np.testing.assert_allclose(jax.jit(silu_ffi)(_x), _x64 * _s, rtol=1e-5, atol=1e-6, err_msg="前向值不是 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="梯度不对")
    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) 结果不对")
    print("✅ 练习 3 自检通过")
⏳ 练习 3 尚未作答:name 'silu_ffi' is not defined
# ✍️ 作答区:练习 4
# 在这里写你的代码,然后运行下面的「🧪 自检」。参考答案在本章末尾的「附录:练习参考答案」。
# 🧪 自检:练习 4
try:
    rms_norm_prim_p
except NameError as e:
    print(f"⏳ 练习 4 尚未作答:{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 结果不对")
    np.testing.assert_allclose(jax.jit(_rp)(_x), rms_norm_np64(_x), rtol=1e-5, atol=1e-6, err_msg="jit 结果不对")
    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 没有作为参数传给内核(换一个 eps 结果应随之改变)")
    assert "stablehlo.custom_call @rms_norm" in jax.jit(_rp).lower(_x).as_text(), "lowering 应调用 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}) 结果不对:注意内核总是在最后一维上归一化")
    print("✅ 练习 4 自检通过")
⏳ 练习 4 尚未作答:name 'rms_norm_prim_p' is not defined

测验

Q1. 在 JAX 0.12 中,下面哪一种写法能零拷贝地把 JAX 数组 x 交给 NumPy? A. np.array(x) B. np.from_dlpack(jax.dlpack.to_dlpack(x)) C. np.from_dlpack(x) D. x.to_numpy()

答案与解析

C。A 会复制;B 里的 to_dlpack 已经移除;D 不存在。np.from_dlpack(x) 通过 DLPack 协议拿到同一块内存(只读),5.2 里用指针比较验证过。np.asarray(x) 在 CPU 上也是零拷贝。

Q2. jax.dlpack.from_dlpack(np_arr) 在什么情况下会悄悄复制?如何确保不复制?

答案与解析

NumPy 数组的起始地址没有按 64 字节对齐时(比如切片 a[1:]),默认的 copy=None 会悄悄复制。传 copy=False 就能保证不复制:做不到零拷贝时,它直接抛 ValueError。

Q3. 下面三种回调中,哪一种的调用可能在 jit 中被当作死代码删除?为什么? jax.debug.callback / io_callback / pure_callback

答案与解析

pure_callback。JAX 相信它是纯函数,输出没被用到就可以删掉,也可以重复执行。另外两种被当作有副作用,总会执行。所以 pure_callback 里千万不能放日志或计数。

Q4. jax.vmap(lambda v: jax.pure_callback(f, sds, v))(xb) 在 0.12 中会发生什么?怎样修?

答案与解析

抛出 NotImplementedError,因为 vmap_method 默认是 None。要显式指定,比如 vmap_method="expand_dims"(回调支持广播时,一次调用处理整批)或 "sequential"(逐个样本调用)。旧的 vectorized=True 已经失效,还会被转发给回调,导致运行时报错。

Q5. 一个自定义原语的 JVP 规则写成 tangent_out = 2 * x * x_dot。要让 jax.grad 工作,还必须注册转置规则吗?

答案与解析

不是必须的。反向模式会转置 JVP 里的线性计算,这里的线性部分只用了内置的 mul,它自带转置规则。只有当切线计算用到了这个原语本身(比如 custom_scale(x_dot, c))时,才需要 ad.primitive_transposes[p]。5.6 和 5.9 分别验证了这两种情况。

Q6. 为什么 custom_scale_p.bind(x, factor=2.0) 无法对 factor 求导?

答案与解析

关键字参数是原语的静态参数(params)。它们会写进 jaxpr 方程、参与缓存,不是被追踪的值,所以没有切线。要对 factor 求导,就得把它当作操作数传:bind(x, factor)。

Q7. 用 XLA FFI 调用 C++ handler Attr<float>("eps") 时写了 ffi_call(...)(x, eps=1e-5),会发生什么?

答案与解析

Python 的 float 会被编码成 F64 属性,到执行时报 Wrong scalar data type: expected F32 but got F64。要传 np.float32(1e-5),5.10 里验证过。

Q8. 接 C++ 代码时,原型阶段可以用 ctypes + pure_callback,现在更推荐 jax.ffi。说出 FFI 的两个关键优势。

答案与解析

① 调用不经过 Python:编译成 StableHLO 的 custom_call,由 XLA 直接调用 C++ 函数指针,没有 GIL,也不用在主机上走一个来回,还能在 GPU 上运行(handler 里能拿到 stream)。② handler 签名是类型安全的:XLA 会检查输入、输出和属性的 dtype,形状信息(dimensions())也直接传给 C++,不用像 ctypes 那样手动传 size。


本章更新说明

下面汇总本章在 JAX 0.12 下需要注意的写法和结论,每条都在正文对应位置用 assert 验证过。

编号小节要点位置
15.1数组所在设备用 .device 属性读取,不能写成 .device();分片数组用 .devices()5.1 ❌/✅ 单元
25.1随机数 key 用 jax.random.key(0) 生成(类型化 key),不再用 jax.random.PRNGKey(0)(入门见 入门 01 JAX简介,深入见 进阶 10 随机数与分片)5.1 ❌/✅ 单元
35.1jnp.array/asarray 总是复制;CPU 上 jax.device_put 对 64B 对齐的数组直接别名(may_alias=False 也一样),修改源数组会改变 JAX 数组;np.asarray(jax_arr) 是零拷贝只读视图5.1 注意 + 验证单元
45.2jax.dlpack.to_dlpack 已移除;改用 DLPack 协议 np.from_dlpack(x) / torch.from_dlpack(x)5.2 ❌/✅ 单元
55.2jax.dlpack.from_dlpack 只接受带 __dlpack__ 的对象,传裸胶囊报 TypeError5.2 ❌ 单元
65.2x.device_buffer 已移除;取缓冲区指针用 x.unsafe_buffer_pointer() / x.addressable_data(0)5.2 ❌/✅ 单元
75.2jax.random.normal 没有 device 参数(传了报 TypeError);先生成,再 device_put5.2 ❌ 单元
85.2零拷贝导入的 JAX 数组和源 NumPy 数组共享内存,修改源数组会改变 JAX 数组的值5.2 注意 + 验证单元
95.2NumPy→JAX 零拷贝要求 64 字节对齐,否则悄悄复制;copy=False 可强制检查5.2 ✅ 单元
105.3jax.experimental.host_callback(id_tap/id_print/call)已删除;改用 jax.debug.callback/print、io_callback、pure_callback5.3 注意 + ❌/✅ 单元
115.30.12 CPU 上三种回调收到的是 jax.Array,要先 np.asarray 再交给 NumPy5.3 ✅ 单元
125.4直接对 pure_callback 求 jax.grad 会报 ValueError: Pure callbacks do not support JVP;需要 custom_jvp/custom_vjp 提供导数5.4 ❌/✅ 单元
135.4pure_callback 默认 vmap_method=None,此时 vmap 报 NotImplementedError;必须指定 5 种取值之一5.4 ❌ 单元 + vmap_method 表
145.4三种回调在 jit/vmap/grad/DCE 下的实际行为见对比表(如 pure_callback 不能直接 grad,输出未用时会被删除)5.3 对比表
155.5原语名是 scan、cond、add(_p 只是 Python 变量名)5.5 注意 + 单元
165.5编译流程是 jaxpr → StableHLO(MLIR)→ XLA5.5 单元
175.6Primitive 从 jax.extend.core 导入(jax.core 不再导出它;ShapedArray 仍在 jax.core)5.6 ❌/✅ 单元
185.6ad.Zero.from_value 不存在;零切线用 ad.Zero(jax.typeof(y).to_tangent_aval()) 构造5.6 ❌/✅ 单元
195.6ad.primitive_transposes 注册的是转置规则,不是 VJP;JVP 只用现有原语时,不注册转置也能 grad5.6 注意 + 单元
205.7多输出原语除了返回 ShapedArray 元组,还必须设 p.multiple_results = True,否则 ValueError5.7 注意 + 单元
215.8lowering 用 mlir.register_lowering + mlir.lower_fun / StableHLO 方言(xla_client.register_translation 和 XLA builder 已不可用);GPU 平台名用 'cuda' 而不是 'gpu'5.8 ❌/✅ 单元
225.9Primitive.def_jvp / def_vjp、jax.core.get_aval、ad.Zero.zero_if_zero 都不存在;用 ad.primitive_jvps + ad.primitive_transposes5.9 ❌ 单元
235.9需要求导的量(如 custom_scale 的 factor)要作为操作数 bind(x, factor) 传入,不能写成参数 bind(x, factor=...)5.9 ❌/✅ 单元
245.9原语要支持 vmap 需注册 batching 规则 batching.primitive_batchers[p],逐元素原语可用 batching.defvectorized5.6、5.9 补充
255.10from jax.experimental import host_callback 会因模块已删除报 ImportError,整段脚本在 import 时就失败5.10 ❌ 单元
265.10pure_callback 的入口是 jax.pure_callback(from jax.experimental import pure_callback 报 ImportError)5.10 ❌ 单元
275.10pure_callback(..., vectorized=True) 已失效(vectorized 会被转发给回调导致报错),改用 vmap_method="expand_dims"/"broadcast_all"5.10 ❌/✅ 单元
285.10没开 x64 时,jnp.arange(..., dtype=jnp.float64) 会被截断为 float32(UserWarning)5.10 ❌ 单元
295.10接入 C++ 用 XLA FFI:jax.ffi.register_ffi_target + ffi_call / ffi_lowering(取代 xla_client.ops.CustomCall / xla.register_translation)5.10 FFI 单元、练习 4

本章小结

  1. NumPy 互操作:jnp.array 总是复制;np.asarray(jax_arr) 和 jax.device_put(CPU、对齐数组)可能零拷贝。转回 NumPy 是同步点,只在流程的边界上做。
  2. DLPack:0.12 只剩协议式接口,即 np.from_dlpack(x)、jax.dlpack.from_dlpack(obj, copy=...),以及数组自带的 __dlpack__ / __dlpack_device__。零拷贝要求同一设备、64 字节对齐。共享内存意味着改源数组会“穿透”到 JAX 数组。
  3. 回调三件套取代了 host_callback:debug.callback 用于观察,io_callback 用于副作用(ordered=True 保证顺序),pure_callback 用于纯计算(必须指定 vmap_method,可能被 DCE)。三者都不能自动微分,要靠 custom_jvp/custom_vjp 补上导数。
  4. 原语是一切变换的最小单元。自定义原语需要:def_impl、def_abstract_eval(返回 ShapedArray)、mlir.register_lowering、ad.primitive_jvps,按需加 ad.primitive_transposes 和 batching.primitive_batchers。这些都是半内部 API。
  5. C++ 集成:原型阶段用 ctypes + pure_callback;生产环境用 XLA FFI:C++ handler → jax.ffi.include_dir() 编译 → pycapsule + register_ffi_target → ffi_call(vmap_method=...) → custom_vjp 接梯度。

速查表

API(JAX 0.12)作用
x.device / x.devices()数组所在设备(属性 / 集合)
jax.device_put(np_arr, device)放到指定设备(CPU 上可能别名 NumPy 内存)
np.asarray(x) / np.array(x)JAX → NumPy:零拷贝只读视图 / 复制
jax.device_get(tree)整个 pytree 拉回主机
np.from_dlpack(x)JAX → NumPy 零拷贝(DLPack 协议)
jax.dlpack.from_dlpack(obj, device=None, copy=None)外部数组 → JAX;copy=False 强制零拷贝
x.__dlpack__() / x.__dlpack_device__()DLPack 协议的生产方接口
x.unsafe_buffer_pointer()底层缓冲区地址(验证零拷贝用)
jax.debug.print(fmt, **kw) / jax.debug.callback(fn, *args)调试打印 / 观察中间值(不返回值)
jax.experimental.io_callback(fn, result_sds, *args, ordered=)有副作用的主机回调
jax.pure_callback(fn, result_sds, *args, vmap_method=)纯主机回调;vmap_method ∈ {sequential, sequential_unrolled, expand_dims, broadcast_all, legacy_vectorized}
jax.effects_barrier()等待所有回调执行完
jax.ShapeDtypeStruct(shape, dtype)描述回调 / FFI 的输出
jax.extend.core.Primitive(name)创建原语(半内部)
p.def_impl / p.def_abstract_evaleager 实现 / 抽象求值(返回 jax.core.ShapedArray)
p.multiple_results = True多输出原语
mlir.register_lowering(p, rule, platform=)注册 lowering(jax.interpreters.mlir)
mlir.lower_fun(fn, multiple_results=)用 jnp 实现生成 lowering
jax.extend.mlir.dialects.stablehlo手写 StableHLO op
ad.primitive_jvps[p] / ad.primitive_transposes[p]JVP / 转置规则(jax.interpreters.ad)
ad.Zero(aval) / ad.is_undefined_primal(x)符号零切线 / 判断转置中的线性输入
batching.primitive_batchers[p] / batching.defvectorized(p)vmap 规则 / 逐元素原语一键 batching
xla.apply_primitive用 lowering 实现 eager impl
jax.eval_shape(fn, *sds)只做抽象求值(调试形状)
jax.test_util.check_grads(f, args, order, modes)数值校验求导规则
jax.ffi.include_dir()XLA FFI C++ 头文件目录
jax.ffi.pycapsule(fn_ptr) / jax.ffi.register_ffi_target(name, cap, platform=)包装并注册 C++ handler
jax.ffi.ffi_call(name, result_sds, vmap_method=)(*args, **attrs)调用 FFI target(属性 dtype 需匹配)
jax.ffi.ffi_lowering(name)把 FFI target 用作原语的 lowering
XLA_FFI_DEFINE_HANDLER_SYMBOL / ffi::Ffi::Bind().Attr<>().Arg<>().Ret<>()C++ 端声明 handler 签名

附录:练习参考答案

先在正文的「✍️ 作答区」自己写,用「🧪 自检」检查;自检通过或者实在卡住了,再来这里对照。 下面的答案按顺序运行,会覆盖你在作答区定义的同名函数。

练习 1

题目见正文「动手练习」。

# 参考答案 1:custom_reduce_sum 的 JVP / 转置 / batching
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)          # 线性:切线也用本原语

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),)   # 广播回输入形状

def _reduce_sum_batch(args, dims, *, axis):
    x, = args; bd, = dims
    x = jnp.moveaxis(x, bd, 0)                                  # 批维度移到 0
    new_axis = tuple(a + 1 for a in axis)                       # 原来的轴整体右移 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("练习 1 通过")
练习 1 通过

练习 2

题目见正文「动手练习」。

# 参考答案 2:SciPy gammaln / digamma 经 pure_callback 接入,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 ufunc 支持任意形状

@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("练习 2 通过")
练习 2 通过

练习 3

题目见正文「动手练习」。

# 参考答案 3(C++ 部分):SiLU 前向 / 反向 FFI 内核,写入 cpp/silu.cc
(CPP_DIR / "silu.cc").write_text(r'''// 练习 3:SiLU 前向 / 反向 FFI 内核(逐元素)
#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("已写入", short_path(CPP_DIR / "silu.cc"))
已写入 cpp/silu.cc
# 参考答案 3:编译、注册 SiLU 内核,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),                                          # 残差就是 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("练习 3 通过")
$ /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
编译成功 → libsilu.so (70824 字节)
练习 3 通过

练习 4

题目见正文「动手练习」。

# 参考答案 4:FFI target 作为自定义原语的 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 变成 FFI 属性
rms_norm_prim_p.def_impl(partial(_xla.apply_primitive, rms_norm_prim_p))                   # eager:临时 jit 这个原语

def _rms_prim_batch(args, dims, *, eps):
    x, = args; bd, = dims
    # 内核把除最后一维外的所有维当作“行”:批维度只要不是最后一维即可,这里统一移到 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("练习 4 通过")
{ lambda ; a:f32[4,32]. let b:f32[4,32] = rms_norm_prim[eps=1e-05] a in (b,) }
练习 4 通过