第 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 之后,先把相关单元跑一遍再往下看。
学习目标¶
- 弄清
jax.Array与numpy.ndarray互转时,哪些操作复制、哪些共享内存、哪里是同步点。 - 用 DLPack 协议(
__dlpack__/np.from_dlpack/jax.dlpack.from_dlpack)实现零拷贝,并用指针比较证明内存确实共享。 - 掌握
host_callback被删除之后的三种替代:jax.debug.callback、io_callback、pure_callback,以及它们在 jit/vmap/grad 下的不同行为。 - 理解原语(Primitive)是 JAX 变换的最小单元,能从零定义一个原语,并依次补上 impl、抽象求值、lowering、JVP、转置和 batching 规则,让它通过 jit、grad 和 vmap。
- 用 XLA FFI(
jax.ffi) 把真正编译出来的 C++ 函数接进 JAX,并补上梯度和 vmap 支持。
小节目录¶
| 小节 | 主题 | 原文 slug |
|---|---|---|
| 5.1 | JAX 与 NumPy 的集成 | integrating-jax-numpy |
| 5.2 | 使用 DLPack 实现零拷贝数据共享 | zero-copy-dlpack |
| 5.3 | calling-external-code-host-callback | |
| 5.4 | 使用 jax.pure_callback 进行无副作用调用 | using-jax-pure-callback |
| 5.5 | JAX 原语简介 | 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 |
之后依次是:动手练习、测验、本章更新说明、本章小结、速查表。
与入门课的关联¶
- 异步调度与
block_until_ready:../01_JAX简介.ipynb、../02_JIT编译.ipynb - jaxpr、追踪(tracing)与 tracer:../02_JIT编译.ipynb。jaxpr 里的每一行是怎么来的,5.5~5.7 讲原语和抽象求值时会说清楚。
- JVP/VJP、
custom_jvp/custom_vjp:../03_自动微分grad.ipynb。5.4、5.9 和 5.10 都靠它们给“黑盒”补梯度。 - vmap 的批处理维度:../04_自动向量化vmap.ipynb。本章的
vmap_method和 batching 规则,讲的就是 vmap 底下的机制。 - 总目录:../00_目录与导读.ipynb
⚠️ 环境注意: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
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 条最佳实践
- 内层循环里尽量少转换,
@jax.jit函数内部更是不要转换。 - 只在计算流程的边界转换:入口用
jnp.array/jax.device_put转入,出口用np.array转出。 - 多设备场景下用
jax.device_put明确表达放置意图。 np.array(jax_array)会悄悄等计算完成。如果只想同步、不想拷数据(比如精确计时),直接调用.block_until_ready()。- 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]
# ❌ 旧写法(过时/错误):jax_arr.device() —— 0.12 中 .device 是属性(Device 对象),不能调用
try:
print(jax_arr.device())
except TypeError as e:
print("TypeError:", e)
# ✅ 推荐写法:.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
# ❌ 旧写法(过时/错误):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)
# ✅ 推荐写法: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)
⚠️ 注意:
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
# 同步点的代价:每步都转回 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))
🔍 补充¶
🟢 必读
- 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())
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 条注意事项
- 设备兼容:双方用的是同一个物理设备,才能零拷贝。DLPack 没法直接在 CPU 和 GPU 之间共享,数据得先搬到目标设备上。
- 数据同步:JAX 是异步执行的,导出前先
block_until_ready()。从别的框架(比如 CUDA 上的 PyTorch)导入前,也要确认对方的计算流已经同步,不然可能读到还没写完的数据。 - 生命周期与所有权:内存一般归原来的框架所有。导入的对象还在用的时候,原张量必须活着,不然会出现悬空指针。
- 可变性:JAX 数组是不可变的。如果接收方原地改了底层缓冲区,JAX 的这个前提就被打破了。
- 流顺序(高级):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,这两段示例在这里跑不了)")
# ✅ 推荐写法: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)
# ✅ 推荐写法: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
⚠️ 注意:零拷贝导入得到的 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
# 计时:零拷贝 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")
🔍 补充¶
⚪ 选读:要和 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
# ✅ 推荐写法: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)
# ✅ 推荐写法: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])
# ✅ 推荐写法: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)
# ✅ 推荐写法:有副作用的调用 → 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
三种回调在 jit / vmap / grad / DCE 下的行为(逐项实测)¶
下一个单元把每种回调放进各种变换里逐个试一遍,得到下面这张表,并用 assert 把结论固定下来。“输出未被使用时”这一列说的是 DCE(死代码消除):回调的输出没被用到时,jit 会不会把这次调用当成死代码删掉。
| 回调 | jit | vmap | grad | 输出未被使用时 |
|---|---|---|---|---|
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 # 纯回调被当作死代码删除
🔍 补充¶
🟢 必读
- 怎么选:只想看一眼中间值,用
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 里并不是这样:
- vmap:
vmap_method默认是None,这时 vmap 直接抛NotImplementedError,必须显式指定。0.12 的合法取值有"sequential"、"sequential_unrolled"、"expand_dims"、"broadcast_all"和"legacy_vectorized"。- grad:不管回调里的数学可不可微,都会抛
ValueError: Pure callbacks do not support JVP。JAX 看不到 NumPy 内部在算什么,必须用jax.custom_jvp或jax.custom_vjp手动提供导数。- 旧参数
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)
# ❌ 旧写法(过时/错误):不指定 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])
# ✅ 推荐写法: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)
# 课程多输出示例: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
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])
🔍 补充¶
🟢 必读
- 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))
# 验证:形状不匹配会在运行时报错;输出未被使用时回调不会执行
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
# 计时: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")
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。
为什么重要(课程讲了三点):
- 变换作用在原语上:
jit把函数追踪成 jaxpr,再交给 XLA 编译;grad靠的是每个原语的求导规则;vmap靠的是每个原语的批处理规则。变换从来不看 Python 代码,只看原语。 - 编译接口:原语连着高层 API 和编译器。jaxpr 被降低(lower)成编译器 IR,再编译成 CPU/GPU/TPU 的机器码。
- 可扩展性:想接入 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")
# 原语身上挂着各种规则表: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
🔍 补充¶
⚪ 选读:要从零定义自定义原语时
- 一个原语要完整工作,需要这些规则(后面几节逐个实现):
| 规则 | 注册位置(0.12) | 服务的变换 | 缺失时的报错 |
|---|---|---|---|
| impl(eager 实现) | p.def_impl | 直接调用(不在 jit 下) | Evaluation rule ... not implemented |
| 抽象求值 | p.def_abstract_eval | 一切追踪:jit / eval_shape / grad / vmap | Abstract evaluation ... not implemented |
| lowering | mlir.register_lowering | jit(编译) | MLIR translation rule ... not found |
| JVP | ad.primitive_jvps[p] | jvp / grad | Differentiation rule ... not implemented |
| 转置 | ad.primitive_transposes[p] | grad / vjp(只有线性原语需要) | Transpose rule ... not implemented |
| batching | batching.primitive_batchers[p] | vmap | Batching 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 / 转置规则)。
课程的五个步骤:
- 创建原语对象,名字会显示在 jaxpr 中。
- 写一个给用户调用的包装函数,里面调用
p.bind(*args, **params)。bind不会马上算出结果,而是交给当前的追踪器(trace)处理:eager 时调用 impl,jit 时记进 jaxpr。会影响输出结构的静态参数(比如some_parameter)要用关键字参数传给bind。 - 用
def_abstract_eval注册抽象求值规则(5.7)。 - 用
mlir.register_lowering注册后端 lowering 规则(5.8)。 - 注册求导规则(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"))
# ✅ 推荐写法:步骤 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]}")
# ✅ 推荐写法:步骤 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])
# ❌ 旧写法(过时/错误):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)
# ✅ 推荐写法:步骤 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) 通过")
⚠️ 注意:
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 就用(抽象输入, 参数)去调用它,得到抽象输出。课程给了三个理由:
- 静态分析:还没编译、没执行,就能发现形状不匹配和类型错误。
- 变换兼容:vmap 要靠它推出批处理后的形状,grad 要靠它搭出反向传播的结构。
- 编译: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
# ✅ 课程示例 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))
⚠️ 注意:多输出的原语,光让抽象求值返回
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)
🔍 补充¶
⚪ 选读:要给自定义原语写抽象求值规则时
- 抽象求值里只能有形状逻辑:它收到的是
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")
# ✅ 推荐写法 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
# ✅ 推荐写法 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])
🔍 补充¶
⚪ 选读:要手写 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])
# ✅ 推荐写法: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])
# ✅ 转置规则:对线性参数 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 通过!")
🔍 补充: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))
🔍 补充¶
⚪ 选读:要给自定义原语写求导规则时
- 为什么只要 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$$
课程步骤:
- C++ 实现:用
extern "C"防止 C++ 的名称修饰(name mangling),这样 ctypes 才能按名字找到符号;对double数组逐元素计算input[i]*input[i] + 10.0。 - 编译为共享库:Linux 用
g++ -shared -fPIC -o custom_op.so custom_op.cpp,macOS 用g++ -shared -o custom_op.dylib custom_op.cpp,Windows 用 MinGW 或 MSVC 的对应命令。 - ctypes 封装:用
ctypes.CDLL加载库,设置argtypes/restype。NumPy 包装函数负责转成 float64、保证内存连续、分配输出,再把指针传给 C 函数。预期结果是[1,2,3] → [11,14,19]。 pure_callback集成:这个 C++ 函数在数学上是纯的,适合用 pure_callback。课程在 eager 和 jit 下测试(预期[1,2,3,4] → [11,14,19,26]),并指出 jax.grad 会失败,因为自动微分看不透 C++ 代码。- 另一种方案(只做概述):定义 Primitive → 抽象求值 → 各后端 lowering(提到
xla_client.ops.CustomCall和xla.register_translation)→ JVP/转置 → 用户函数bind。 - 总结:回调适合纯函数、并且不需要穿过外部代码求梯度的场景。自定义原语集成得最紧密,也支持自动微分,但复杂度高。
本节先按课程的思路,用 0.12 的写法把 ctypes 方案跑通。然后用现在推荐的 XLA FFI(jax.ffi),接入一个真正编译出来的 C++ RMSNorm 内核,并补上梯度和 vmap。
⚠️ 注意:这段集成代码如果照旧写法来写,在 0.12 上有 4 个地方会出问题:
from jax.experimental import host_callback as hcb直接报ImportError。哪怕导入之后根本没用到它,整段脚本也会在 import 阶段就失败;from jax.experimental import pure_callback报ImportError,正确的入口是jax.pure_callback;pure_callback(..., vectorized=True):vectorized已经不是 pure_callback 的参数了,它会被当成 kwargs 转发给回调,运行时报错。应该改成vmap_method="expand_dims"或"broadcast_all";- 默认没开 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;
}
}
}
# 课程步骤 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()
# 课程步骤 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)
# ❌ 旧写法(过时/错误):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
# ✅ 推荐写法: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))
现代做法: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 的流程:
- 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++。 - 编译:头文件在
jax.ffi.include_dir()下面。只要头文件,不用链接任何库。 - Python 端:
ctypes.cdll.LoadLibrary加载 →jax.ffi.pycapsule(lib.符号)包装函数指针 →jax.ffi.register_ffi_target(名字, 胶囊, platform="cpu")注册。 - 调用:
jax.ffi.ffi_call(名字, 输出 ShapeDtypeStruct, vmap_method=...)(*输入, **属性)。它生成的是 StableHLO 的custom_call,天然支持 jit,也不会像回调那样在 Python 里走一个来回。 - 梯度:
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>>());
# 编译:头文件路径来自 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_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
# 验证两个易错点:属性 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
# ✅ 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) 通过")
# ✅ 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 内核 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")
🔍 补充¶
⚪ 选读:要把 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 自检通过")
# ✍️ 作答区:练习 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 自检通过")
# ✍️ 作答区:练习 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 自检通过")
# ✍️ 作答区:练习 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 自检通过")
测验¶
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 验证过。
| 编号 | 小节 | 要点 | 位置 |
|---|---|---|---|
| 1 | 5.1 | 数组所在设备用 .device 属性读取,不能写成 .device();分片数组用 .devices() | 5.1 ❌/✅ 单元 |
| 2 | 5.1 | 随机数 key 用 jax.random.key(0) 生成(类型化 key),不再用 jax.random.PRNGKey(0)(入门见 入门 01 JAX简介,深入见 进阶 10 随机数与分片) | 5.1 ❌/✅ 单元 |
| 3 | 5.1 | jnp.array/asarray 总是复制;CPU 上 jax.device_put 对 64B 对齐的数组直接别名(may_alias=False 也一样),修改源数组会改变 JAX 数组;np.asarray(jax_arr) 是零拷贝只读视图 | 5.1 注意 + 验证单元 |
| 4 | 5.2 | jax.dlpack.to_dlpack 已移除;改用 DLPack 协议 np.from_dlpack(x) / torch.from_dlpack(x) | 5.2 ❌/✅ 单元 |
| 5 | 5.2 | jax.dlpack.from_dlpack 只接受带 __dlpack__ 的对象,传裸胶囊报 TypeError | 5.2 ❌ 单元 |
| 6 | 5.2 | x.device_buffer 已移除;取缓冲区指针用 x.unsafe_buffer_pointer() / x.addressable_data(0) | 5.2 ❌/✅ 单元 |
| 7 | 5.2 | jax.random.normal 没有 device 参数(传了报 TypeError);先生成,再 device_put | 5.2 ❌ 单元 |
| 8 | 5.2 | 零拷贝导入的 JAX 数组和源 NumPy 数组共享内存,修改源数组会改变 JAX 数组的值 | 5.2 注意 + 验证单元 |
| 9 | 5.2 | NumPy→JAX 零拷贝要求 64 字节对齐,否则悄悄复制;copy=False 可强制检查 | 5.2 ✅ 单元 |
| 10 | 5.3 | jax.experimental.host_callback(id_tap/id_print/call)已删除;改用 jax.debug.callback/print、io_callback、pure_callback | 5.3 注意 + ❌/✅ 单元 |
| 11 | 5.3 | 0.12 CPU 上三种回调收到的是 jax.Array,要先 np.asarray 再交给 NumPy | 5.3 ✅ 单元 |
| 12 | 5.4 | 直接对 pure_callback 求 jax.grad 会报 ValueError: Pure callbacks do not support JVP;需要 custom_jvp/custom_vjp 提供导数 | 5.4 ❌/✅ 单元 |
| 13 | 5.4 | pure_callback 默认 vmap_method=None,此时 vmap 报 NotImplementedError;必须指定 5 种取值之一 | 5.4 ❌ 单元 + vmap_method 表 |
| 14 | 5.4 | 三种回调在 jit/vmap/grad/DCE 下的实际行为见对比表(如 pure_callback 不能直接 grad,输出未用时会被删除) | 5.3 对比表 |
| 15 | 5.5 | 原语名是 scan、cond、add(_p 只是 Python 变量名) | 5.5 注意 + 单元 |
| 16 | 5.5 | 编译流程是 jaxpr → StableHLO(MLIR)→ XLA | 5.5 单元 |
| 17 | 5.6 | Primitive 从 jax.extend.core 导入(jax.core 不再导出它;ShapedArray 仍在 jax.core) | 5.6 ❌/✅ 单元 |
| 18 | 5.6 | ad.Zero.from_value 不存在;零切线用 ad.Zero(jax.typeof(y).to_tangent_aval()) 构造 | 5.6 ❌/✅ 单元 |
| 19 | 5.6 | ad.primitive_transposes 注册的是转置规则,不是 VJP;JVP 只用现有原语时,不注册转置也能 grad | 5.6 注意 + 单元 |
| 20 | 5.7 | 多输出原语除了返回 ShapedArray 元组,还必须设 p.multiple_results = True,否则 ValueError | 5.7 注意 + 单元 |
| 21 | 5.8 | lowering 用 mlir.register_lowering + mlir.lower_fun / StableHLO 方言(xla_client.register_translation 和 XLA builder 已不可用);GPU 平台名用 'cuda' 而不是 'gpu' | 5.8 ❌/✅ 单元 |
| 22 | 5.9 | Primitive.def_jvp / def_vjp、jax.core.get_aval、ad.Zero.zero_if_zero 都不存在;用 ad.primitive_jvps + ad.primitive_transposes | 5.9 ❌ 单元 |
| 23 | 5.9 | 需要求导的量(如 custom_scale 的 factor)要作为操作数 bind(x, factor) 传入,不能写成参数 bind(x, factor=...) | 5.9 ❌/✅ 单元 |
| 24 | 5.9 | 原语要支持 vmap 需注册 batching 规则 batching.primitive_batchers[p],逐元素原语可用 batching.defvectorized | 5.6、5.9 补充 |
| 25 | 5.10 | from jax.experimental import host_callback 会因模块已删除报 ImportError,整段脚本在 import 时就失败 | 5.10 ❌ 单元 |
| 26 | 5.10 | pure_callback 的入口是 jax.pure_callback(from jax.experimental import pure_callback 报 ImportError) | 5.10 ❌ 单元 |
| 27 | 5.10 | pure_callback(..., vectorized=True) 已失效(vectorized 会被转发给回调导致报错),改用 vmap_method="expand_dims"/"broadcast_all" | 5.10 ❌/✅ 单元 |
| 28 | 5.10 | 没开 x64 时,jnp.arange(..., dtype=jnp.float64) 会被截断为 float32(UserWarning) | 5.10 ❌ 单元 |
| 29 | 5.10 | 接入 C++ 用 XLA FFI:jax.ffi.register_ffi_target + ffi_call / ffi_lowering(取代 xla_client.ops.CustomCall / xla.register_translation) | 5.10 FFI 单元、练习 4 |
本章小结¶
- NumPy 互操作:
jnp.array总是复制;np.asarray(jax_arr)和jax.device_put(CPU、对齐数组)可能零拷贝。转回 NumPy 是同步点,只在流程的边界上做。 - DLPack:0.12 只剩协议式接口,即
np.from_dlpack(x)、jax.dlpack.from_dlpack(obj, copy=...),以及数组自带的__dlpack__/__dlpack_device__。零拷贝要求同一设备、64 字节对齐。共享内存意味着改源数组会“穿透”到 JAX 数组。 - 回调三件套取代了
host_callback:debug.callback用于观察,io_callback用于副作用(ordered=True保证顺序),pure_callback用于纯计算(必须指定vmap_method,可能被 DCE)。三者都不能自动微分,要靠custom_jvp/custom_vjp补上导数。 - 原语是一切变换的最小单元。自定义原语需要:
def_impl、def_abstract_eval(返回ShapedArray)、mlir.register_lowering、ad.primitive_jvps,按需加ad.primitive_transposes和batching.primitive_batchers。这些都是半内部 API。 - 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_eval | eager 实现 / 抽象求值(返回 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 通过")
练习 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 通过")
练习 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"))
# 参考答案 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 通过")
练习 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 通过")