eager 和延迟执行
主线位置：张量主线的「计算」一步：调用一个 op 之后，结果什么时候真正算出来，出了错在哪一行知道。
版本：PyTorch 2.15.0（CPU、MPS）、burn 0.22.1（Flex 设备、Metal 设备）
flowchart TB
callop["调用 a.exp()"] --> cpu{"哪种设备"}
cpu -->|"torch CPU 或 Flex"| now["当场算完再返回"]
cpu -->|torch MPS| queue["内核进队列 立即返回"]
cpu -->|burn Metal| rec["只登记 IR 立即返回"]
queue --> gpu["GPU 按顺序执行"]
rec --> wait["等读回或 sync"]
wait --> fuse["融合后执行"]
wait -.->|"结果已被丢掉"| skip["整块不执行"]
PyTorch 的 eager 有两层意思：每个 op 按调用顺序各自执行一次，而且调用时就把这次执行交出去。burn 在 CPU 上也是这样；在 GPU 上，Fusion 只满足「按顺序」不满足「调用时交出去」：op 先记成 IR，攒到必须出结果时才决定怎么合并、怎么执行。本页量这个差别在四个地方的表现：调用返回的时间、没人读的结果、一串逐元素 op 的耗时、错误出现的位置。
实验代码：PyTorch 侧 exec_gpu.py；burn 侧 crate exec-gpu，依赖 burn = { version = "=0.22.1", features = ["flex", "metal"] }。这里打开了 burn 的默认 feature，因为 fusion 和 autotune 是靠默认 feature 打开的，本卷模板那种 default-features = false 的 Metal 没有 Fusion（Fusion 装饰器对比过两种构建）。全部 release 构建，平台 M2 Max、macOS 27、rustc 1.99.0。测量时机器上有别的编译任务，uptime 的 1 分钟负载在 18 到 30 之间；PyTorch 的线程数是默认的 8。
先确认 burn 两个设备各走哪几层：
for d in ["flex", "metal"] {
let dev = device(d);
println!("{d}: {}", Dispatch::name(dev.as_dispatch()));
}
flex: dispatch<flex>
metal: dispatch<fusion<cubecl<wgpu<msl>>>>
调用返回的时候，算完了没有
同一行 a.exp()，a 是 [4096, 4096] 的 f32。先算一次把第一次的开销（内核编译）放到计时之外，再连做 5 次，每次记两个时间：调用返回时、同步之后。
def ret(dev: str, n: int) -> None:
a = rand(n, dev)
a.exp().cpu()
sync(dev)
for i in range(5):
t = time.perf_counter()
y = a.exp()
returned = time.perf_counter() - t
sync(dev) # torch.mps.synchronize() on MPS, nothing on CPU
synced = time.perf_counter() - t
del y
print(f"{dev} {n}x{n} exp run {i}: returned after {fmt(returned)}, after sync {fmt(synced)}")
fn ret(dev: &str, n: usize) {
let device = device(dev);
warmup(&device);
let a = rand(n, &device);
let _ = a.clone().exp().into_data(); // first call (kernel compile) out of the way
device.sync().unwrap();
for i in 0..5 {
let t = Instant::now();
let y = a.clone().exp();
let returned = t.elapsed();
device.sync().unwrap();
let synced = t.elapsed();
drop(y);
println!("{dev} {n}x{n} exp run {i}: returned after {returned:?}, after sync {synced:?}");
}
}
cpu 4096x4096 exp run 0: returned after 5.943ms, after sync 5.947ms
cpu 4096x4096 exp run 1: returned after 8.038ms, after sync 8.042ms
...
mps 4096x4096 exp run 0: returned after 50.000µs, after sync 913.458µs
mps 4096x4096 exp run 1: returned after 78.625µs, after sync 1.357ms
mps 4096x4096 exp run 2: returned after 32.042µs, after sync 829.167µs
mps 4096x4096 exp run 3: returned after 17.250µs, after sync 821.000µs
mps 4096x4096 exp run 4: returned after 16.833µs, after sync 1.218ms
flex 4096x4096 exp run 0: returned after 28.9325ms, after sync 28.932833ms
flex 4096x4096 exp run 1: returned after 27.787333ms, after sync 27.787625ms
...
metal 4096x4096 exp run 0: returned after 10.209µs, after sync 1.446125ms
metal 4096x4096 exp run 1: returned after 7.625µs, after sync 3.121ms
metal 4096x4096 exp run 2: returned after 7.208µs, after sync 1.602ms
metal 4096x4096 exp run 3: returned after 6.833µs, after sync 2.027167ms
metal 4096x4096 exp run 4: returned after 6.5µs, after sync 1.248209ms
- CPU 上两边一样：调用返回时已经算完，
sync不再花时间。Flex 是即时执行的后端，Device::flush的文档写着「Eager backends, which execute each operation as it is registered, have nothing buffered and treat this as a no-op」。 - GPU 上两边都是调用先返回、同步时才等到结果。只看这一组数，分不出「排进了 GPU 的队列」和「只记了下来」，下面两节才分得出来。
两边 CPU 上的速度差（Flex 的 exp 比 PyTorch 慢几倍）不是本页的话题，速度对比和它的前提见计时。
没人用的结果
(a * b + c).exp() 算 20 次。一种写法每次的结果直接丢掉、最后同步一次；另一种每次留住结果、同步、再丢。
for _ in range(reps):
expr() # result dropped right away
sync(dev)
dropped = (time.perf_counter() - t) / reps
for _ in 0..reps {
let _ = expr(); // dropped at the end of the statement
}
device.sync().unwrap();
let dropped = t.elapsed() / reps;
mps 4096x4096 (a*b+c).exp() round 0: dropped 1.706ms per eval; kept+sync each min 1.994ms median 2.070ms
mps 4096x4096 (a*b+c).exp() round 1: dropped 1.787ms per eval; kept+sync each min 1.947ms median 5.315ms
mps 4096x4096 (a*b+c).exp() round 2: dropped 1.583ms per eval; kept+sync each min 1.832ms median 2.037ms
metal 4096x4096 (a*b+c).exp() round 0: dropped 360.454µs per eval; kept+sync each min 987.708µs median 1.082ms
metal 4096x4096 (a*b+c).exp() round 1: dropped 7.406µs per eval; kept+sync each min 987.958µs median 1.071ms
metal 4096x4096 (a*b+c).exp() round 2: dropped 6.412µs per eval; kept+sync each min 977.042µs median 1.048ms
PyTorch 在 MPS 上丢掉的结果照样花了 1.6 到 1.8 毫秒一次，和留住结果的差不多：调用时内核已经交给了 GPU，之后没人读它也要跑完。burn 在 Metal 上从第 1 轮起每次只要 7 微秒左右，这是登记 IR 的开销，计算本身没有发生：结果在执行前就被 drop，Fusion 把整块省掉（机制见 Fusion 装饰器「没人读的结果」）。第 0 轮的 360 微秒里多出来的部分本页没有追。CPU 上两边都即时执行，丢不丢结果时间一样（PyTorch CPU 每次 15 毫秒左右，Flex 30 到 44 毫秒，输出在实验目录里）。
这条差别会直接影响从 PyTorch 搬过来的基准代码：for _ in range(n): f(x) 加一次同步，在 MPS 上量的是 n 次计算，搬到 burn 的 Metal 上量的是 n 次登记。写法见计时和 burn 卷的基准测试测出来几乎为零。
问： PyTorch 在 MPS 上调用也是立刻返回，那它和 burn 的延迟执行有什么区别？
答： 区别在调用返回时工作交出去了没有。PyTorch 返回时已经分配好输出、把这个 op 交给了 GPU，所以丢掉的结果照样计算（上面每次 1.6 毫秒），显存不够在调用那一行就报错（最后一节），每个 op 都是一次独立的执行（下一节）。burn 返回时只有一条 IR 在通道里，执行与否、和谁合并，都等到有人要结果时才定。前者是异步（asynchronous），后者是延迟（lazy）。
一串逐元素 op
8 个逐元素 op 连着做，中间结果都只用一次：
def chain8(x: torch.Tensor) -> torch.Tensor:
y = (x * 0.5 + 0.25).sin().abs()
return (y * 0.5 + 0.25).sin().abs()
fn chain8(x: Tensor<2>) -> Tensor<2> {
let y = (x * 0.5 + 0.25).sin().abs();
(y * 0.5 + 0.25).sin().abs()
}
PyTorch 侧用 torch.mps.Event(enable_timing=True) 量设备上的时间，burn 侧用 device.profile_with(ProfileOptions::default().flush(), ..)，各 10 次（两种设备时间的用法见计时）。torch.compile 那一行是同一个函数包一层 torch.compile(chain8)，第一次调用前设了 TORCHINDUCTOR_CACHE_DIR，缓存放在实验目录里。
mps 4096x4096 chain8 first call (read back): 56.449ms
mps 4096x4096 chain8 wall (sync), 10 runs: min 3.569ms median 4.035ms
mps 4096x4096 chain8 mps.Event time, 10 runs: min 3.384ms median 3.532ms
...
mps 4096x4096 chain8 compiled first call (read back): 346.258ms
mps 4096x4096 chain8 compiled wall (sync), 10 runs: min 1.478ms median 1.782ms
mps 4096x4096 chain8 compiled mps.Event time, 10 runs: min 1.063ms median 1.474ms
metal 4096x4096 chain8 first call (read back): 70.967ms
metal 4096x4096 chain8 wall (sync), 10 runs: min 738.583µs median 1.093ms
metal 4096x4096 chain8 device time (profile_with flush), 10 runs: min 412.958µs median 421.750µs
eager 的 PyTorch 每个 op 写出一个完整的 [4096, 4096] 中间张量，下一个 op 再读回来，8 个 op 共 8 次读写 64 MiB；设备时间的中位数 3.5 毫秒（另一次运行是 3.0 毫秒）。torch.compile 生成的代码（缓存目录里的 Python 文件）只有一个 Metal 内核，8 步在寄存器里做完：
kernel void generated_kernel_0(
device float* out_ptr0,
constant float* in_ptr0,
uint xindex [[thread_position_in_grid]]
) {
int x0 = xindex;
auto tmp0 = in_ptr0[x0];
auto tmp1 = 0.5;
auto tmp2 = tmp0 * tmp1;
auto tmp3 = 0.25;
auto tmp4 = tmp2 + tmp3;
auto tmp5 = metal::precise::sin(tmp4);
auto tmp6 = metal::abs(tmp5);
...
out_ptr0[x0] = static_cast<float>(tmp10);
}
burn 在 Metal 上不用改代码就是这个形状：Fusion 把 8 个 op 合成一个内核（同样的 8 个 op 在 burn 卷里数过启动次数，开 Fusion 1 次、不开 8 次，见 Fusion 装饰器「开和不开」），设备时间约 0.41 毫秒。torch.compile 版本的设备时间中位数在几次运行里是 1.0 到 1.5 毫秒，两个融合版本之间还差两到三倍，本页没有追原因；能看到的一点是 inductor 的内核每个线程算一个元素、用 metal::precise::sin，burn 的融合内核按末维向量化（生成融合内核）。
快慢之外先确认结果一样。用确定的输入 x = 10·sin(k)（[256, 256]），各设备的结果和 float64 参考值比：
torch cpu max|x - f64 ref| = 1.247e-07
torch mps max|x - f64 ref| = 1.726e-07
torch mps compiled max|x - f64 ref| = 1.726e-07
burn flex max|x - f64 ref| = 2.560e-07
burn metal max|x - f64 ref| = 4.251e-07
结果都在 [0, 1] 之间，误差都在 f32 几个 ulp 以内；不是 0 是因为各家 sin 的实现不同。
CPU 这一侧：Flex 没有 Fusion，chain8 每次 59 到 63 毫秒；PyTorch CPU eager 10.5 到 37 毫秒（负载变化大），torch.compile 之后 5.2 毫秒。在 burn 里，「要不要融合」不是代码里的开关，而是由设备决定：同一个函数，换到 Metal 上自动融合，在 Flex 上逐个执行。PyTorch 反过来，融合要自己加 torch.compile，第一次调用要编译（MPS 上几次运行在 0.35 到 0.94 秒之间，取决于 inductor 缓存里已有什么；CPU 上 1.9 秒）。
错误在哪一行出现
三种错误，看它们在哪一行被报出来。
形状错误：[2, 3] @ [2, 3]。
--- mps: a @ b with [2, 3] x [2, 3]
--- raised at the call: RuntimeError: mm: mat1 and mat2 shapes cannot be multiplied (2x3 and 2x3)
--- metal: a.matmul(b) with [2, 3] x [2, 3]
thread 'main' (8626761) panicked at .../burn-tensor-0.22.1/src/tensor/api/numeric.rs:842:9:
=== Tensor Operation Error ===
Operation: 'Matmul'
Reason:
1. The inner dimension of matmul should be the same, but got 3 and 2. Lhs shape Shape { dims: [2, 3] }, rhs shape Shape { dims: [2, 3] }.
note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace
--- returned to caller: panicked
四种设备都在调用那一行报。burn 的形状检查在 Tensor::matmul 的第一行 check!(TensorCheck::matmul(..))，在进入任何后端之前，所以 Metal 上也不延迟。形状是登记时就知道的，只有执行时才知道的错误才会延迟。
显存不够：建一个 [2^18, 2^18] 的 f32（256 GiB）再加 1。
huge = torch.zeros(2**18, 2**18, device="mps") + 1
--- torch.zeros(2**18, 2**18, device='mps') + 1
--- raised at the call: RuntimeError: Invalid buffer size: 256.00 GiB
let huge = Tensor::<2>::zeros([1 << 18, 1 << 18], &device) + 1.0;
println!("--- huge registered, dims {:?}", huge.dims());
let small = Tensor::<1>::from_floats([1.0, 2.0], &device) * 2.0;
println!("--- small.try_to_data(): {:?}", small.try_to_data().map(|d| d.to_vec::<f32>().unwrap()));
println!("--- device.sync(): {:?}", device.sync());
match huge.slice([0..1, 0..2]).try_to_data() {
Ok(d) => println!("--- huge read: Ok {:?}", d.to_vec::<f32>()),
Err(e) => println!("--- huge read: Err: {e}"),
}
--- huge registered, dims [262144, 262144]
--- small.try_to_data(): Ok([2.0, 4.0])
--- device.sync(): Ok(())
--- huge read: Err: An error happened during execution
Caused by:
Several failures at once
...
Caused by:
can't allocate buffer of size: 274877906944
PyTorch 在调用时就分配输出，分配失败当场报。burn 的分配发生在执行时，执行又在读别的张量时才发生；错误挂在 huge 上，读 small 和 sync 都是 Ok，读 huge 才拿到。机制和这里 sync 返回 Ok 与文档不符的问题见 burn 卷的出错的 op 不报错，后面读别的张量才报。
下标越界：x = [1, 2, 3]，用下标 5 去取，先读另一个不相干的张量 z = x * 2，再读结果（PyTorch 侧 exec_oob.py，burn 侧 exec-gpu 的 oob 模式，各跑两遍，结果相同）。
x = torch.tensor([1., 2., 3.], device="mps")
i = torch.tensor([5], device="mps")
y = x[i]
z = x * 2
z.cpu() # then z.cpu() again, then y.cpu()
x[i] returned
z.cpu() raised: AcceleratorError index 5 is out of bounds for dimension 0 with size 3
z.cpu() again -> tensor([2., 4., 6.])
y.cpu() -> tensor([0.])
let x = Tensor::<1>::from_floats([1.0, 2.0, 3.0], &device);
let i = Tensor::<1, Int>::from_ints([5], &device);
let y = x.clone().select(0, i);
let z = x * 2.0;
// z.try_to_data(), device.sync(), y.try_to_data()
--- metal: x.select(0, [5]) returned
--- z.try_to_data(): Ok([2.0, 4.0, 6.0])
--- device.sync(): Ok(())
--- y.try_to_data(): Ok [0.0]
CPU 上两边都在调用那一行报（PyTorch IndexError: index 5 is out of bounds for dimension 0 with size 3；Flex panic 在 burn-flex-0.22.1/src/ops/gather_scatter.rs:144:5，信息 index 5 out of bounds for dimension of size 3）。GPU 上两边都不在调用处报，但去向不同：
- PyTorch MPS 把错误报给下一个同步点，不管同步的是谁：读不相干的
z时抛出AcceleratorError，报过一次就清掉了，再读z正常，读y得到[0.]。 - burn Metal 根本不报：
select的 cubecl 内核select_kernel用launch_unchecked启动，读input[offset_input]前不检查下标，越界读到的是那块内存里碰巧有的东西，这里两次都是 0（cubecl 的越界读写行为见 burn 卷自定义内核越界读不报错）。select的文档没有提越界。
从 PyTorch 过来，调试时的习惯要换一下：在 CPU 上复现再看哪一行出错，这一点两边一样；在 GPU 上，PyTorch 的报错位置不可信但至少会报，burn 的显存错误挂在张量上、要读那个张量才知道，下标越界则要靠自己在 CPU 上（Flex）验证。
chain8 在 Metal 上，如果在第 4 个 op 之后插一句 println!("{:?}", y.dims())，融合会不会被打断，变成两个内核？提示：dims() 要的是形状，形状在登记时就算好了。平台和构建同上。
动脑筋答案
不会。exec-gpu 的 chaindims 模式在同一个进程里量三个版本：原版、中间插 println!("{:?}", y.dims())、中间插 black_box(y.to_data())，各 10 次设备时间，整个程序跑 3 遍（负载约 2.5）：
fn chain8_dims(x: Tensor<2>) -> Tensor<2> {
let y = (x * 0.5 + 0.25).sin().abs();
println!("{:?}", y.dims());
(y * 0.5 + 0.25).sin().abs()
}
metal 4096x4096 chain8 device time (profile_with flush), 10 runs: min 409.875µs median 410.750µs
metal 4096x4096 chain8+dims device time (profile_with flush), 10 runs: min 346.666µs median 409.583µs
metal 4096x4096 chain8+to_data device time (profile_with flush), 10 runs: min 633.375µs median 699.667µs
...
metal 4096x4096 chain8 device time (profile_with flush), 10 runs: min 409.958µs median 415.792µs
metal 4096x4096 chain8+dims device time (profile_with flush), 10 runs: min 345.917µs median 410.292µs
metal 4096x4096 chain8+to_data device time (profile_with flush), 10 runs: min 642.709µs median 701.125µs
- 插
dims()的中位数和原版都是 0.41 ms：dims()读的是登记时由 IR 算出的形状，不碰数据，不触发执行。（两者的最小值有时落在 0.35 ms，原版在别的进程里也出现过 0.35 ms，本页没有追这个双峰。） - 插
to_data()的版本 0.70 ms：读数据让前 4 个 op 先单独执行、结果拷回主机，后 4 个 op 另成一段。触发执行的操作（读数据、sync、跨线程使用张量等）完整列表见 Fusion 装饰器「什么时候必须执行」。 - PyTorch 里对应的是
print(y.shape)不同步、print(y)同步：读形状两边都不碰设备，读数值两边都要等。区别是 burn 的「等」还会切断融合，调试时在中间打印数值，量到的速度就不是原来的速度了。
小结
- PyTorch 在 CPU 上和 burn 的 Flex 一样，调用返回时已经算完；PyTorch 在 MPS 上是异步：调用时分配输出、把 op 交给 GPU，之后一定执行。
- burn 在 Metal（带 Fusion）上是延迟：调用只登记 IR，执行时才决定合并；没人读的结果整块不执行，量基准时要留住结果。
- 8 个逐元素 op 的设备时间：PyTorch MPS eager 3.0–3.5 ms，
torch.compile后 1.0–1.5 ms，burn Metal 不改代码约 0.41 ms；融合在 PyTorch 里是torch.compile的选项，在 burn 里由设备决定。 - 形状错误两边在任何设备上都当场报；显存不够 PyTorch 当场报，burn 读到那个张量才报；GPU 上的下标越界，PyTorch 报给下一个同步点，burn Metal 不报，读到的值不可信（实测是 0）。