设计取舍
主线位置:三条主线的汇总。算子主线上的取舍是「后端是运行期的值」「编译代价在哪里付」「装饰器与通道」「惰性加融合」「自己写 GPU 语言」「JIT 与 autotune」;求导主线上的是「autodiff 是设备上的标记」;训练主线继承前两条线的全部代价。每一条只说收益和代价,机制在链接指向的页面里。
flowchart TB user["Tensor D K 无后端参数"] -->|每 op 几十 ns| disp["Dispatch 按设备 match"] disp -->|按后端数编译| comp["burn-dispatch 单态化"] disp --> ad["Autodiff 设备标记"] disp --> flex["Flex 即时执行"] ad --> fus["Fusion 惰性登记"] disp --> fus fus -->|错误和计时推迟| cube["cubecl 内核语言"] cube -->|首次编译与调优| tune["JIT 与 autotune"]
每一节回答三个问题:burn 选了什么,换来了什么,付出了什么。依据只有三种:本卷各页和本页跑出来的数,burn / cubecl / cubek 在 tag 上的源码和提交记录,candle 0.11.0 tag 上的源码。0.21 不在离线缓存里,所以「0.21 写法编译多久、报什么错」没有实测,只引用源码和提交记录。实验平台 M2 Max、macOS 27、rustc 1.99.0,GPU 只测了 Metal(经 wgpu);CUDA、ROCm、Vulkan 只读源码。
1. 后端是运行期的值,不是类型参数
选择。 0.22 的 Tensor<D, K> 没有后端参数,后端由张量所在的 Device 决定,所有调用经过一个全局后端 Dispatch,它在每个 op 里按张量的种类标签 match 到具体后端(Dispatch)。0.21 是 Tensor<B, D, K>,模型写成 Mlp<B: Backend>,后端在编译期就定了。
换来。 模型和训练代码里没有 <B>;同一个二进制在运行期选设备,Device::default() 按一个固定的顺序挑已开启的后端,BURN_DEVICE 环境变量可以覆盖(Device);训练态和推理态是同一个类型(模块与训练)。
付出。
- 每个 op 一次
match加一次拆包装包。16 个元素的 f32 加法(本身约 140 ns,release),Dispatch层多 19–25 ns(只开flex)到 36–40 ns(开flex metal autodiff),Tensor层再多 45–93 ns;DispatchTensor按值搬进搬出,从 192 字节涨到 576 字节(Dispatch)。对 1024² 的矩阵乘可以忽略,对一长串小张量的 op 是同一量级。GPU 上这几十纳秒被惰性登记的几微秒盖住(第 5 节)。 - 后端和设备的错误从编译期挪到运行期。两个张量不在同一设备上,0.22 是一个 panic:
The provided tensors are not on the same device;Device::default()在开了metal时是 wgpu 的Auto设备,和Device::metal(..)不相等,也触发这个 panic(检查与报错、Device)。cat不查设备,跨后端时在分派宏里 panic,同后端跨设备时在后台线程里出错(检查与报错)。仍在编译期的是秩和种类:秩不对是E0308,对Bool张量调matmul是E0599(Tensor 与 Kind)。
问: 既然要在运行期选后端,为什么不用 Box<dyn Backend>,而要为每个 op 生成一个 match?
答: Backend 不能做成 trait 对象。它的定义要求 Clone + Default + Sized,op 都是不带 self 的关联函数(fn float_add(lhs: FloatTensor<B>, rhs: FloatTensor<B>) -> FloatTensor<B>),张量类型是关联类型,每个后端都不一样。分派只能落在张量值上:DispatchTensorKind 每个后端一个变体,而且不装箱,源码注释写着「Boxing backend tensors would cost every tensor an allocation」。代价就是上面那个 576 字节。
2. 编译代价在哪里付
选择。 用户 crate 不碰后端的类型树:Tensor 的泛型方法只做检查和形状计算,调用外提到只经手 BridgeTensor 的非泛型函数,BridgeTensor 和 Device 用 obfuscate! 存成固定大小的槽(BridgeTensor 与 *_impl)。类型树只在 burn-dispatch 里单态化一次:#[backend_dispatch] 给每个 op、每个开启的后端生成 5 份代码(普通一份,autodiff 两种策略各两份)(Dispatch)。
换来。 用户 crate 的编译不随后端变重。BridgeTensor 与 *_impl 用一个直接持有 DispatchTensor 的替身做了对照:release 下替身让叶子 crate 的函数定义从 48 个变成 450 个,其中 155 个是析构胶水,编译慢 2 到 4 倍。0.21 的泛型写法里,模型代码对每个后端类型各单态化一遍,这一部分本页没有实测。
付出。 这部分代价没有消失,挪到了 burn-dispatch,而且按开启的后端数付,不管程序用不用。实验 design-feat / design-lean:同一个只用 Device::flex() 的程序,只改 burn 的 feature,release 构建,cargo build --timings 读每个 crate 的编译时间:
use burn::tensor::{Device, Tensor};
fn main() {
// Only Flex is ever used; the extra features just add backends to Dispatch.
let device = Device::flex();
let a = Tensor::<2>::ones([64, 64], &device);
let b = a.clone().matmul(a) + 1.0;
let first = b.into_data().to_vec::<f32>().unwrap()[0];
println!(
"metal={} autodiff={} first={first} size_of::<Tensor<2>>={}",
cfg!(feature = "metal"),
cfg!(feature = "autodiff"),
std::mem::size_of::<Tensor<2>>(),
);
}
Device::flex() 的程序,在默认 feature 之外多开一个 metal,release 二进制会变大多少?为什么链接器删不掉没用到的 Metal 代码?先猜一个倍数,再看文末。design-lean 用 default-features = false, features = ["std", "flex"],再按需加 autodiff、metal;design-feat 用 burn 的默认 feature 加 flex(默认的 optim 已经带上 autodiff),再加 metal(默认 feature 让 cubecl 后端带上 fusion 和 autotune)。为了只重编 burn-dispatch,两个 crate 都用 [patch.crates-io] 指向一份原样拷贝的 burn-dispatch 0.22.1 源码,每次 touch 它的 lib.rs 再构建,连跑 3 次。原始输出(每组摘最后一次运行的程序输出):
lean-flex run 1: burn-dispatch 21.0 s, rebuild after touch 42 s
lean-flex run 2: burn-dispatch 12.1 s, rebuild after touch 35 s
lean-flex run 3: burn-dispatch 17.5 s, rebuild after touch 36 s
lean-flex binary 1545136 bytes, rlib 1747800 bytes
metal=false autodiff=false first=65 size_of::<Tensor<2>>=256
lean-flex-autodiff run 1: burn-dispatch 49.2 s, rebuild after touch 52 s
lean-flex-autodiff run 2: burn-dispatch 53.9 s, rebuild after touch 55 s
lean-flex-autodiff run 3: burn-dispatch 28.4 s, rebuild after touch 30 s
lean-flex-autodiff binary 3258336 bytes, rlib 11644656 bytes
...
lean-flex-metal run 3: burn-dispatch 41.5 s, rebuild after touch 43 s
lean-flex-metal binary 65879120 bytes, rlib 9102176 bytes
...
lean-flex-metal-autodiff run 3: burn-dispatch 62.9 s, rebuild after touch 65 s
lean-flex-metal-autodiff binary 68769872 bytes, rlib 28181256 bytes
...
default-flex-metal run 1: burn-dispatch 44.6 s, rebuild after touch 47 s
default-flex-metal binary 105342064 bytes, rlib 42749064 bytes
| feature(都只用 Flex) | burn-dispatch 编译(3 次最短) | libburn_dispatch rlib | 二进制(release,未 strip) |
|---|---|---|---|
flex | 12 s | 1.7 MB | 1.5 MB |
flex autodiff | 28 s | 11.6 MB | 3.3 MB |
flex metal | 42 s | 9.1 MB | 65.9 MB |
flex metal autodiff | 63 s | 28.2 MB | 68.8 MB |
默认 feature + flex(含 autodiff) | — | — | 6.7 MB |
默认 feature + flex metal(含 autodiff、fusion、autotune) | 45 s | 42.7 MB | 105.3 MB |
- 编译时间的绝对值不可靠:测量时机器上同时跑着别的构建,
uptime的负载在 70 到 108 之间,同一组 3 次相差到两倍。rlib 大小是确定的,它更能说明问题:开autodiff让burn-dispatch变成约 6.7 倍,和 Dispatch 里数的「每个后端从 1 份变 5 份」是同一个量级,加一个Cube后端约 5 倍,两者都开约 16 倍,再加上 fusion 是 24 倍。 - 程序一行 Metal 代码都没有,二进制却从 6.7 MB 变成 105 MB(strip 之后 4.2 MB 对 66.5 MB)。
nm数符号:名字里带cubecl的从 451 个变成 103 397 个,带wgpu的从 0 个变成 11 484 个。Dispatch的每个match都有Cube分支,链接器没法证明它走不到,wgpu、cubecl 运行时、cubek 的内核就都留在二进制里。0.21 里只写Flex的程序不会实例化Wgpu的代码;0.22 里「编进哪些后端」只由 feature 决定,和代码用了什么无关。 - 编译这一笔在每次改 feature、升级 burn 时付一次,改自己的代码时不付,那一侧见 BridgeTensor 与 *_impl 的实验。二进制大小这一笔每次发布都付。
3. autodiff 是设备上的标记
选择。 0.21 里 backward 只定义在 B: AutodiffBackend 的张量上,0.22 改成设备上的一个标记,backward() 在运行期检查(burn 解决什么问题的冷知识讲了这次改动的来龙去脉)。
换来。 一个模型类型同时用于训练和推理,valid() / train() 在同一个类型上切换,不必在 Autodiff<B> 和 B 两个类型之间转换模块(Module 派生)。
付出。 类型系统不再替你挡的事,变成运行期的 panic 或者静默的行为:
- 在普通设备上
require_grad()、对没被跟踪的张量backward()都是 panic,信息里写着该怎么改(Tensor::require_grad requires autodiff; call Tensor::autodiff first,Autodiff 装饰器)。 - autodiff 张量和普通张量在同一硬件上可以直接运算,普通张量被当成常量抬进图里,不报错(Autodiff 装饰器)。
- 在 autodiff 设备上只做前向,图一直留着:Metal 上 20 次前向从 1 MiB 涨到 41 MiB,
without_autodiff()之后不涨(backward 与 Gradients)。0.21 的AutodiffModule::valid返回Self::InnerModule,一个落在B::InnerBackend上的另一个类型,推理代码换了类型也就离开了 autodiff;0.22 合并了这个 trait(#5721),valid()返回同一个类型,切不切要自己记得。 - 已初始化的模型
to_device到另一个 autodiff 设备,梯度为 0 而不报错,要用fork(Module 派生)。
4. 装饰器与通道
选择。 自动微分和融合是以后端为类型参数的装饰器:Autodiff<B, C> 和 Fusion<B> 都实现 Backend,每个 op 做完自己的事再调 B:: 的同名 op(burn 解决什么问题)。远程执行走另一条路:BackendRouter<R: RouterChannel> 的参数是一个通道,op 被记成 burn-ir 送进通道,由另一头的真实后端重放(burn-ir)。
换来。 自动微分写一遍,对 Flex、cubecl、Remote 都成立;融合写一遍,对所有 cubecl 运行时都成立,开了 fusion 时连远程后端也套上:RemoteBackend 是 Fusion<BackendRouter<RemoteChannel>>,文档说这样「sending a repeated computation (e.g. a model block per step) over the network once instead of every step」。
付出。
- 0.21 里装饰器怎么套由用户在类型里写;0.22 里能套成什么样由分派表写死。
BACKENDS里只有Cube、Flex、NdArray、LibTorch、Remote、Capture六项,Cube就是Fusion<CubeBackend>(burn-cubecl 后端),Autodiff是DispatchTensorKind上包在外面的一层,Capture不能包(burn-dispatch的注释:「Capture does not support autodiff」)。想加一种新的套法,要改burn-backend-extension的这张表,而不是在自己的代码里写一个类型。 - Flex 不实现
FusionBackend,CPU 上没有融合(burn 解决什么问题、Flex)。 - 每一层都转发一次,GPU 上一次 matmul 穿过
dispatch<fusion<cubecl<wgpu<msl>>>>,加 autodiff 再多一层(一次 matmul)。
5. 惰性加融合,而不是即时执行
选择。 cubecl 后端默认套着 Fusion:op 只登记一条 OperationIr 就返回,等到读数据、sync 或者融合器自己收口时,流才被 drain,能合的 op 合成一个内核(Fusion 装饰器、融合搜索)。Flex 是即时的,调用返回时已经算完。
换来。 显存流量和内核数少了。4096² f32 上连续 8 个逐元素 op,8 个内核变 1 个,稳定后每次 3.03 ms 变 505 µs(Fusion 装饰器);(a * b + c).exp() 3 个内核变 1 个,1.61 ms 变 883 µs,流量减半、时间接近减半(生成融合内核)。第一次运行也更快(19.5 ms 对 39.7 ms),因为要编译的内核少。调用方这一侧,Metal 上 matmul 那一行 3–12 µs 就返回(一次 matmul)。
付出。
- 错误推迟:出错的块在设备线程上执行,错误挂在它写出的张量上,读到这些张量时才交出,同一条流上不相干的张量照常可读(Fusion 装饰器)。
- 计时变难:
matmul那一行的时间只是登记;没人读的结果根本不算,基准测试如果丢掉结果,测到的几乎是零;互不相干的计算会被合进同一个内核(Fusion 装饰器、融合搜索)。 Device::profile的窗口里如果只有惰性工作,读到None(Device)。- 融合是默认 feature 带进来的:
default-features = false时 cubecl 后端的fusion和autotune一起关掉,不报错(怎么读这一卷),上面那些数就回到不融合的一栏。
6. 自己写 GPU 语言,而不是绑定 CUDA 或写 WGSL 字符串
选择。 burn 的 GPU 内核都是 cubecl 的 #[cube] 函数:宏把 Rust 函数体展开成构造 IR 的代码,运行期在设备线程上按 KernelId 编译(#[cube] 内核)。cubecl 0.11.3 里,同一份 IR 能编成六种目标:cubecl-cpp 出 CUDA、HIP、MSL 三种 C++ 方言,cubecl-wgpu 出 WGSL,cubecl-spirv 出 SPIR-V,CPU 运行时经 cubecl-llvm 交给 LLVM。
换来。
- 一份内核源码对所有 cubecl 运行时成立,
burn-cubecl只有一个后端类型(burn-cubecl 后端)。 - 融合离不开它:所有逐元素融合都是同一个
#[cube]函数elemwise_fuse,op 列表作为 comptime 配置在编译时展开成一串寄存器运算(生成融合内核)。手写的内核文件做不到「运行期拿到一串 op 再生成内核」。 - autotune 有东西可挑:matmul 的 33 个候选是 cubek 各种策略和参数的组合,都是
#[cube]内核,每个运行时都能编(autotune)。
付出。
- 自己维护一整套语言和编译器:
cubecl-core约 4.0 万行、cubecl-ir1.8 万、cubecl-macros1.1 万(CubeCL 与 GPU)。编译器的问题也要自己背:0.11.3 的 MSL 路径上,带边界检查的launch和launch_unchecked编出同样的代码(#[cube] 内核)。 - API 变动大:0.11 把
Line换成Vector<F, N>、全局缓冲换成&[T]、ComputeClient改名Client,网上和旧版本的内核写法都不再能编(CubeCL 与 GPU)。 - 可移植不等于一样快:同一块 GPU 只开
wgpu时走 WGSL,没有 CMMA,2048² matmul 慢三倍多(wgpu 与 Metal)。
v0.12.1 的 burn-wgpu/src/template/ 下有 37 个 .wgsl 模板,元素类型和工作组大小是 {{ elem }}、{{ workgroup_size_x }} 这样的占位符,运行时替换成文本再交给 wgpu;同一个版本里融合用的是另一条路,burn-wgpu/src/codegen/ 从 IR 生成着色器。2024 年 2、3 月的一串 PR 把模板逐个迁到生成那条路上,标题多是「Migrate/jit/...」,例如 #1378「Migrate wgsl index shaders to gpu representation」;#1417「Splitted the JIT stuff from the Wgpu stuff」把生成那部分拆成独立于 wgpu 的 burn-jit;#1541 的提交说明里有一条「delete template folder」。首次发布在 v0.13.0,那个版本只剩示例 custom-wgpu-kernel 里的一个 .wgsl。再往后,#2041「Migration/cubecl」(首次发布在 v0.14.0)删掉了 burn-compute 和 burn-wgpu 里的 WGSL 编译器,burn 从此依赖单独的 cubecl 仓库。这些提交的说明都没有写原因;从 diff 能看到的是,迁完之后逐元素内核和融合内核走的是同一条生成路径,这正是第 5 节的融合能套到所有运行时上的前提。
和手写内核的方案比,candle 0.11.0 是一个具体的例子:CUDA 内核是 .cu 文件,由 candle-kernels/build.rs 在构建时用 nvcc 编译(大部分编成 PTX,MoE 相关的一组编进静态库 libmoe.a),构建机上要有 CUDA 工具链;Metal 内核是 .metal 源码,include_str! 进二进制,第一次用到时在 load_library 里调 newLibraryWithSource 编译,进程内缓存;macOS 自己还会跨进程缓存编译好的库,命中时每个库只要几毫秒,没命中时 gemm、reduce 这样的库各要 4 秒左右(见 candle 卷的 Metal 后端)。用户写自定义算子实现 CustomOp1,cpu_fwd、cuda_fwd、metal_fwd 各写一遍。Metal 这一步和 burn 一样是运行期编译 MSL 文本,差别在 MSL 是谁写的。
7. JIT 与 autotune 的启动成本
选择。 内核在第一次用到时编译,matmul 这类有多个变体的 op 在第一次遇到某个调优键时实测挑最快的,结果写进当前目录往上最外层 Cargo.toml 旁的 target/environment/default.db(autotune)。
换来。 不用为每块 GPU 预先编好、调好;调优结果按设备和形状桶记下来,之后的进程直接用。
付出。 第一次很慢,而且慢多少取决于缓存状态,release 构建下 1024² f32 matmul(autotune):
| 情形 | 第一次 | 之后每次(最短) |
|---|---|---|
| 空缓存:先测设备峰值(7.5 万次启动)、10 个候选上场、编 11 个内核 | 3.67 s | 2.0 ms |
| 同一进程,同桶同键(768²) | 117 ms | 1.1 ms |
| 同一进程,同桶新键(1000²) | 1.43 s | 2.0 ms |
| 新进程,缓存已有 | 150 ms | 2.0 ms |
- 编译产物在
wgpu<msl>上不跨进程缓存,compilation.cache = true只对 SPIR-V 有用,所以每个新进程的第一次至少要编一次 MSL,100–160 ms(wgpu 与 Metal)。 - 缓存位置跟着当前目录走,同一个二进制在不同目录下启动,看到的是不同的缓存;调试日志会改变测出的启动开销,进而改变调优结果(autotune)。
- 不经 autotune 的逐元素内核也要编译:4096² 上
(a * b + c).exp()的第一次是 19.5 ms,稳定后 883 µs(生成融合内核)。
Flex 没有这一层,第一次和之后一样。candle 0.11.0 的 candle-core、candle-kernels、candle-metal-kernels 里搜不到 autotune,内核是固定的,Metal 库的编译发生在每个进程第一次用到时,系统的库缓存命中时这一步只要几毫秒,没命中时要几秒。
8. 和 candle 0.11 对照
burn 解决什么问题已经列了 candle 的设备和存储 enum、内建在 Tensor 里的自动微分。上面几节补了内核和启动成本,这里再补两点和运行期行为有关的事实。
- 错误是
Result。 candle 的Tensor::matmul返回Result<Self>,设备不一致时Storage::matmul返回Err(DeviceMismatchBinaryOp)。burn 的Tensor方法返回张量本身,同样的错误是 panic;Result只出现在读数据、同步和设备设置上(检查与报错)。 - 即时编码,批量提交。 candle 的 Metal 后端在调用时就把内核编码进当前的命令缓冲,多个 op 共用一个计算编码器,每取用编码器
DEFAULT_CANDLE_METAL_COMPUTE_PER_BUFFER= 50 次(大致一个 op 一次)提交一次命令缓冲(CANDLE_METAL_COMPUTE_PER_BUFFER可改)。GPU 执行同样是异步的,但每个 op 都是一个内核,没有合并;burn 推迟的不只是提交,还有「执行哪些内核」这个决定(第 5 节)。
burn-candle 在 #650「Feat/candle/initialize」(2023-08,首次发布在 v0.9.0)加入,让 candle 的张量作为 burn 的一个 Backend 实现,套在 Autodiff 下面用。#4416「chore: deprecate burn-candle backend」(首次发布在 v0.21.0-pre.1)把它标为弃用,建议改用 cubecl 后端、ndarray 或 tch。#5343「chore!: remove deprecated burn-candle backend」(首次发布在 v0.22.0-pre.3)删掉了 crate 和 candle、candle-cuda、candle-metal 三个 feature,提交说明给的附带收益是依赖图去重:「gemm 0.18, safetensors 0.4 and zip 7 were reachable only through candle's ug chain」。所以 0.22.1 的 BACKENDS 表里没有 candle;当年它能接进来,靠的正是第 4 节说的「实现 Backend 就能被装饰器套上」。
总表
| 取舍 | 换来的 | 付出的 | 证据 |
|---|---|---|---|
| 1. 后端是运行期的值 | 模型无 <B>;同一二进制运行期选设备;训练推理同一类型 | 每个小 op 多 19–40 ns(Tensor 层再 45–93 ns);设备错误变成运行期 panic | Dispatch、检查与报错 |
2. 编译代价在 burn-dispatch 付 | 用户 crate 不展开类型树(替身对照 48 对 450 个函数) | burn-dispatch 按后端数和 autodiff 变重(rlib 1.7 MB 到 42.7 MB);只用 Flex 的程序开了 metal 也从 6.7 MB 变 105 MB | BridgeTensor 与 *_impl、本页实验 |
| 3. autodiff 是设备标记 | 一个模型类型训练推理两用 | require_grad / backward 运行期 panic;前向留图,内存涨;to_device 后梯度为 0 | Autodiff 装饰器、backward 与 Gradients |
| 4. 装饰器与通道 | 自动微分、融合各写一遍;远程后端也能融合 | 能套成什么样由分派表写死;Flex 没有融合;每层转发一次 | burn 解决什么问题、Dispatch |
| 5. 惰性加融合 | 8 个逐元素 op 1 个内核,3.03 ms 变 505 µs | 错误读时才出;丢掉的结果不算;计时要 sync | Fusion 装饰器、生成融合内核 |
| 6. 自写 GPU 语言 | 一份内核编到六种目标;融合内核靠 comptime 生成;autotune 有候选 | 自己维护编译器和它的问题;API 变动大;WGSL 路径慢三倍多 | #[cube] 内核、wgpu 与 Metal |
| 7. JIT 与 autotune | 不必预编译、预调优;结果跨进程复用 | 空缓存 1024² 第一次 3.67 s;每个新进程至少编一次 MSL | autotune |
| 8. 对照 candle 0.11 | — | candle:enum 后端、Result 错误、手写 .cu / .metal、即时编码、无融合无 autotune | 本页、burn 解决什么问题 |
动脑筋答案
「多大」:6.7 MB 变成 105.3 MB,大约 16 倍(strip 之后 4.2 MB 对 66.5 MB)。「为什么」:Dispatch 的每个 op 都是一个对所有已开启后端的 match,Cube 分支和它背后的 Fusion<CubeBackend>、cubecl 的 wgpu 运行时、cubek 的内核是同一个函数里的代码,链接器不能删。用 nm 数两个二进制的符号:
bin-flex size=6696112 stripped=4224784 symbols= 11645 burn_dispatch=1139 cubecl=451 wgpu=0
bin-metal-autodiff size=105342064 stripped=66547720 symbols= 205749 burn_dispatch=4780 cubecl=103397 wgpu=11484
bin-flex 里那 451 个 cubecl 符号来自 cubecl-common、cubecl-zspace 这些不管后端都会用到的小 crate(Shape 就在 cubecl-zspace 里)。想让只用 CPU 的程序保持小,办法只有一个:不开 GPU 的 feature。
小结
- 0.22 把后端从类型挪到
Device值上:模型代码干净了,同一二进制能在运行期选设备;代价是每个小 op 多几十纳秒,设备和 autodiff 的错误变成运行期 panic,秩和种类的错误仍在编译期。 - 编译代价从用户 crate 挪进
burn-dispatch:BridgeTensor挡住了下游,burn-dispatch自己按开启的后端数变重,二进制里带着用不到的后端。 - Autodiff、Fusion 是类型装饰器,Router、Remote 是通道;0.22 的组合方式写死在分派表里,Flex 没有融合。
- 惰性加融合让逐元素链快数倍,代价是错误、计时和「到底算没算」都推迟到读数据的时候。
- cubecl 让一份内核编到多种目标、支撑运行期生成的融合内核和 autotune 的候选;代价是自己维护编译器,以及冷缓存下第一次几秒、每个新进程一次 MSL 编译的启动成本。
- candle 0.11 是对照面:enum 后端、
Result错误、手写内核、即时编码,没有融合和 autotune。