Rust in Everything
基于 burn 0.22.1

设计取舍

主线位置:三条主线的汇总。算子主线上的取舍是「后端是运行期的值」「编译代价在哪里付」「装饰器与通道」「惰性加融合」「自己写 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)。
📝NOTE
没有蠢问题

问: 既然要在运行期选后端,为什么不用 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>>(),
  );
}
❗IMPORTANT
动脑筋:一个只用 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)
flex12 s1.7 MB1.5 MB
flex autodiff28 s11.6 MB3.3 MB
flex metal42 s9.1 MB65.9 MB
flex metal autodiff63 s28.2 MB68.8 MB
默认 feature + flex(含 autodiff)——6.7 MB
默认 feature + flex metal(含 autodiff、fusion、autotune)45 s42.7 MB105.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-ir 1.8 万、cubecl-macros 1.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)。
📝NOTE
冷知识:burn 自己也写过 WGSL 字符串

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 s2.0 ms
同一进程,同桶同键(768²)117 ms1.1 ms
同一进程,同桶新键(1000²)1.43 s2.0 ms
新进程,缓存已有150 ms2.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 节)。
📝NOTE
冷知识:candle 曾经是 burn 的一个后端

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);设备错误变成运行期 panicDispatch、检查与报错
2. 编译代价在 burn-dispatch 付用户 crate 不展开类型树(替身对照 48 对 450 个函数)burn-dispatch 按后端数和 autodiff 变重(rlib 1.7 MB 到 42.7 MB);只用 Flex 的程序开了 metal 也从 6.7 MB 变 105 MBBridgeTensor 与 *_impl、本页实验
3. autodiff 是设备标记一个模型类型训练推理两用require_grad / backward 运行期 panic;前向留图,内存涨;to_device 后梯度为 0Autodiff 装饰器、backward 与 Gradients
4. 装饰器与通道自动微分、融合各写一遍;远程后端也能融合能套成什么样由分派表写死;Flex 没有融合;每层转发一次burn 解决什么问题、Dispatch
5. 惰性加融合8 个逐元素 op 1 个内核,3.03 ms 变 505 µs错误读时才出;丢掉的结果不算;计时要 syncFusion 装饰器、生成融合内核
6. 自写 GPU 语言一份内核编到六种目标;融合内核靠 comptime 生成;autotune 有候选自己维护编译器和它的问题;API 变动大;WGSL 路径慢三倍多#[cube] 内核、wgpu 与 Metal
7. JIT 与 autotune不必预编译、预调优;结果跨进程复用空缓存 1024² 第一次 3.67 s;每个新进程至少编一次 MSLautotune
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。

💬 关联讨论

发起讨论
加载中...