Rust in Everything
基于 burn 0.22.1

#[cube] 内核

主线位置：算子主线的倒数第二段。cubek 的 matmul、burn-cubecl 的逐元素内核、融合生成的内核，最后都是一个 #[cube] 函数。本页讲这个函数被宏展开成了什么、一次 launch 在调用线程上做了什么、IR 什么时候在哪里构造。之后句柄、内存池和读回见运行时与 Client；IR 怎么编成 MSL、管线怎么缓存见 wgpu 与 Metal。

flowchart TB
  src["cube launch 标注的 fn"] --> orig["原函数 只做类型检查"]
  src --> modn["同名 mod"]
  modn --> exp["expand 构造 IR"]
  modn --> desc["内核描述体"]
  modn --> lau["launch 函数"]
  lau --> reg["登记参数"]
  reg --> cl["Client launch"]
  cl --> dev["设备线程 查 KernelId"]
  dev -->|没编译过| def["define 调 expand"]
  def --> msl["编成 MSL"]

本页的 cubecl 版本是 0.11.3（burn = "=0.22.1" 解析到的版本）。0.11 的前端和半年前差别很大：内核参数从 &Array<Line<F>> 变成了 &[Vector<F, N>]，IR 换成了 pliron，ComputeClient<R> 变成了不带运行时参数的 Client。v0.11.0-pre.1（2026-07-28）之前发布的 cubecl 里内核参数还是 &Array<..> 写法，下面每个 API 都按 tag 核对过。

实验程序

实验 crate cube-kernels 只开 burn 的 metal 和 cubecl 两个 feature。cubecl feature 把 cubecl 转出成 burn::cubecl。#[cube] 宏生成的代码写的是相对路径 cubecl::prelude::...（core_path 只是一个 cubecl 标识符），所以不必单独依赖 cubecl，在 crate 根上 use burn::cubecl; 就能用：

use burn::cubecl;
use cubecl::prelude::*;

/// y = a * x + y, one vector of N floats per unit.
#[cube(launch)]
fn saxpy<F: Float, N: Size>(a: f32, x: &[Vector<F, N>], y: &mut [Vector<F, N>]) {
  if ABSOLUTE_POS < x.len() {
    y[ABSOLUTE_POS] = Vector::new(F::cast_from(a)) * x[ABSOLUTE_POS] + y[ABSOLUTE_POS];
  }
}

fn run_saxpy(client: &cubecl::client::Client) {
  let x: Vec<f32> = (0..16).map(|i| i as f32).collect();
  let y: Vec<f32> = vec![1.0; 16];
  let (x_handle, y_handle) = (upload(client, &x), upload(client, &y));
  let vector_size = 4;
  saxpy::launch::<f32>(
    client,
    CubeCount::Static(1, 1, 1),
    CubeDim::new_1d(x.len() as u32 / vector_size as u32),
    vector_size,
    2.0f32,
    // SAFETY: both handles hold exactly `len` f32 values.
    unsafe { BufferArg::from_raw_parts(x_handle, x.len()) },
    unsafe { BufferArg::from_raw_parts(y_handle.clone(), y.len()) },
  );
  let bytes = client.read_one(y_handle).unwrap();
  println!("saxpy y = {:?}", f32::from_bytes(&bytes));
}

fn main() {
  let device = cubecl::Device::default();
  let client = device.client();
  println!("device = {device:?}, runtime = {:?}", client.name());
  // ...
}

upload 是 client.create_from_slice(f32::as_bytes(data))。16 个 f32 按每 4 个一组，4 个单元（unit）各算一组。debug 构建，M2 Max、macOS 27、rustc 1.99.0：

device = Wgpu(WgpuDevice { kind: DefaultDevice, backend: Auto }), runtime = "wgpu<msl>"
saxpy y = [1.0, 3.0, 5.0, 7.0, 9.0, 11.0, 13.0, 15.0, 17.0, 19.0, 21.0, 23.0, 25.0, 27.0, 29.0, 31.0]

cubecl::Device::default() 在只开了 wgpu 的构建里是 Wgpu(.. backend: Auto)，在 macOS 上落到 MSL 编译器，和 burn 的 Device::metal(..) 走同一条路（见一次 matmul）。几处写法和旧教程不同：标量参数要写具体类型（a: F 会报 the trait bound F: LaunchArg is not satisfied，impl_scalar_launch! 只给具体的数值类型实现了 LaunchArg）；BufferArg::from_raw_parts 是 unsafe fn，因为长度由调用者担保；Vector 的宽度不再挂在缓冲参数上，而是 launch 的第四个参数 vector_size。

宏展开成了什么

#[cube] 处理函数时输出两样东西：去掉辅助属性的原函数，和一个从函数生成的 Launch 结构（quote! { #[allow(dead_code, clippy::too_many_arguments)] #item #kernel }）。cargo expand 这个 crate 得到 846 行，saxpy 一个函数（原函数加生成的模块）约 350 行。按项列出来（省略的部分标 ...）：

#[allow(dead_code, clippy::too_many_arguments)]
/// y = a * x + y, one vector of N floats per unit.
fn saxpy<F: Float, N: Size>(a: f32, x: &[Vector<F, N>], y: &mut [Vector<F, N>]) {
    if ABSOLUTE_POS < x.len() {
        y[ABSOLUTE_POS] = Vector::new(F::cast_from(a)) * x[ABSOLUTE_POS]
            + y[ABSOLUTE_POS];
    }
}
mod saxpy {
    use super::*;
    pub struct _F;
    /// Type to be used as a generic for launch kernel argument.
    pub type F = cubecl::prelude::DynamicScalar<_F>;
    pub struct _N;
    /// Type to be used as a generic for launch kernel argument.
    pub type N = cubecl::prelude::DynamicSize<_N>;
    pub fn expand<F: Float, N: Size>(
        scope: &cubecl::prelude::Scope,
        _n: usize,
        a: <f32 as cubecl::prelude::CubeType>::ExpandType,
        x: &<[Vector<F, N>] as cubecl::prelude::CubeType>::ExpandType,
        y: &mut <[Vector<F, N>] as cubecl::prelude::CubeType>::ExpandType,
    ) -> () { ... }
    ///saxpy Kernel
    pub struct Saxpy<F: Float> {
        settings: cubecl::prelude::KernelSettings,
        device_properties: cubecl::__private::Arc<cubecl::prelude::DeviceProperties>,
        target_properties: cubecl::__private::Arc<cubecl::prelude::TargetProperties>,
        a: <f32 as cubecl::prelude::LaunchArg>::CompilationArg,
        x: <[Vector<F, N>] as cubecl::prelude::LaunchArg>::CompilationArg,
        y: <[Vector<F, N>] as cubecl::prelude::LaunchArg>::CompilationArg,
        _n: usize,
        __ty: ::core::marker::PhantomData<(F)>,
    }
    struct SaxpyInfo<F: Float> { ... }   // 同样几个字段，派生 Hash / Eq / Debug
    impl<F: Float> Saxpy<F> { pub fn new(...) -> Self { ... } }
    impl<F: Float> cubecl::prelude::KernelMetadata for Saxpy<F> {
        fn id(&self) -> cubecl::prelude::KernelId { ... }
        // ...
    }
    impl<F: Float> cubecl::prelude::CubeKernel for Saxpy<F> {
        fn define(&self) -> cubecl::prelude::KernelDefinition { ... }
    }
    pub fn launch<'kernel, F: Float>(
        __client: &cubecl::prelude::Client,
        __cube_count: cubecl::prelude::CubeCount,
        __cube_dim: cubecl::prelude::CubeDim,
        _n: usize,
        a: cubecl::RuntimeArg<f32>,
        x: cubecl::RuntimeArg<[Vector<F, N>]>,
        y: cubecl::RuntimeArg<[Vector<F, N>]>,
    ) { ... }
}
生成的项做什么什么时候运行
原函数 fn saxpy让 rustc 按普通 Rust 检查函数体的类型从不调用
saxpy::expand把函数体改写成一串「往 Scope 里登记 IR 操作」的调用设备线程上，这个内核第一次编译时
Saxpy<F> 和 SaxpyInfo<F>一次启动的「编译参数」：设置、设备属性、每个参数的 CompilationArg、向量宽度 _nlaunch 时构造，交给 Client
KernelMetadata::id把类型、设置和 SaxpyInfo 合成 KernelId，当缓存键设备线程查缓存时
CubeKernel::define建 KernelBuilder，为每个参数造 IR 值，调 expand，build()缓存没命中时
saxpy::launch登记运行期参数，构造 Saxpy，提交调用线程上，每次调用

两个 mod saxpy 里的类型别名是理解泛型的关键。launch::<f32> 的 F 是真类型，用来算元素大小和注册元素类型；define 调 expand 时传的却是 expand::<DynamicScalar<_F>, DynamicSize<_N>>。元素类型和向量宽度先登记进 scope 的类型表（register_type::<DynamicScalar<_F>>(F::elem_type(..))、register_size::<DynamicSize<_N>>(self._n.into())），expand 里再查出来。所以构造 IR 的代码对每个内核只单态化一份，f32、f16 用的是同一段 expand，区别只在类型表里的值。burn 的 examples/custom-cubecl-kernel 更进一步，用 #[define(F)] _dtype: ElemType 让元素类型在启动时由一个运行期值决定，后端按张量的 dtype 传进去。

📝NOTE
没有蠢问题

问： 原函数照常编译，函数体里的 ABSOLUTE_POS 在 CPU 上是什么值？saxpy 会不会被当成 CPU 函数调用？

答： 不会调用，它只为类型检查而存在（所以带 #[allow(dead_code)]）。ABSOLUTE_POS 在 constant_usize! 里定义成 pub const ABSOLUTE_POS: usize = 2;，值是随手写的，只为让 ABSOLUTE_POS < x.len() 这种表达式类型对得上。同一个宏还生成一个同名的模块 ABSOLUTE_POS，里面的 expand 才是真身：登记一个 ReadBuiltinOp，读 Builtin::AbsolutePos。expand 函数里写的是 ABSOLUTE_POS::expand(scope)，use 常量时模块一起进了作用域，文件头的注释管这叫「a trick」。

expand：函数体变成构造 IR 的调用

saxpy::expand 里每个运算符都换成了方法调用，参数是上一步返回的 IR 值。去掉层层 IntoExpand::into_expand 包装，结构是：

cubecl::frontend::branch::if_expand(
    scope,
    ABSOLUTE_POS::expand(scope).__expand_lt_method(scope, &x.__expand_len_method(scope)),
    |scope| {
        let _value = Vector::__expand_new(scope, F::__expand_cast_from(scope, a.into()))
            .__expand_mul_method(scope, __expand_deref(scope, x.__expand_index_method(scope, ABSOLUTE_POS::expand(scope))))
            .__expand_add_method(scope, __expand_deref(scope, y.__expand_index_method(scope, ABSOLUTE_POS::expand(scope))));
        cubecl::prelude::__expand_assign(
            scope,
            y.__expand_index_mut_method(scope, ABSOLUTE_POS::expand(scope)),
            _value,
        )
    },
);

（为了能读，这段是从展开结果里删掉包装后整理的，不是原样输出。）

这些 __expand_* 方法最后都落到同一个模式：在 pliron 的 Context 里建一个操作，插到 Scope 当前的插入点。二元运算是 binary_expand，取下标是 index_expand：

pub(crate) fn binary_expand<F, O>(
    scope: &Scope,
    lhs: ExpandValue,
    rhs: ExpandValue,
    func: F,
) -> ExpandValue
where
    F: Fn(&mut Context, Value, Value) -> O,
    O: Op + OneResultInterface,
{
    let [lhs, rhs] =
        normalize_same_vectorization(scope, [lhs.read_value(scope), rhs.read_value(scope)]);
    let op = func(scope.ctx_mut(), lhs, rhs);
    scope.register_with_result(&op).into()
}

pub(crate) fn index_expand(scope: &Scope, list: Value, index: Value, checked: bool) -> Value {
    let op = IndexOp::maybe_checked(scope.ctx_mut(), list, index, checked);
    scope.register_with_result(&op)
}

Scope 只有三个字段：pliron 上下文的句柄、插入点、处理 return / break 的展开状态；register 就是 self.inserter().append_op(ctx, op)。normalize_same_vectorization 是 Vector 和标量能直接相乘的原因：一边是标量、另一边是宽度 4 的向量时，先把标量广播成向量；两边宽度都大于 1 又不相等就 panic。运行期的 if 变成 if_expand：条件在展开时已经是常量的，直接调用闭包或者跳过；否则建一个 IfOp，给 then、else 两个块各开一个子 Scope，闭包往 then 块里继续登记。

内核的参数从哪来：define 对每个参数调 LaunchArg::expand。切片参数走 impl LaunchArg for [C]，由 Scope::global 给入口函数加一个类型为「指向运行期数组的指针」的参数，再配上从元数据里读出的长度，拼成一个切片值。x.len() 读的就是那个长度。

前端类型

写法是什么在 MSL 里变成
&[T]、&mut [T]全局缓冲，运行期参数是 BufferArg（句柄 + 长度）device T const* / device T*，长度进 info_st.static_meta
&Tensor<T>带形状和 strides 的全局缓冲，参数是 TensorArg同上，另有形状、strides 元数据
Vector<F, N>N 个 F 的向量，N: Size 由 launch 的 vector_size 定float_4 这样的结构体
f32、u32 等标量参数按值传入info_st.scalars_f32[..]
#[comptime] relu: bool编译期常量，进 KernelId不出现，分支已经选好
comptime!(expr)在展开时（CPU 上）算出的值一个字面量
Array<E>内核里的定长局部数组，Array::new(#[comptime] length)局部数组
ABSOLUTE_POS 等内建的位置和尺寸由 threadgroup_position_in_grid 等算出

最后一行之前的 Array 值得说一句：旧教程里的 &Array<Line<F>> 是全局缓冲，0.11 里 Array<E> 只剩内核里的定长数组这一个用途，全局缓冲改用切片。Vector<P, N> 也不是旧的 Line<P> 改个名：宽度从运行期的 line_size 字段变成了类型参数 N。

📝NOTE
冷知识：同一个 gelu 示例，签名在半年里改了两次

cubecl 仓库的 examples/gelu 一直是「最小内核」的范例，它的签名记录了前端的两次重写。#1221（「refactor: Line size generic」，首发 v0.10.0-pre.3）把 fn gelu_array<F: Float>(input: &Array<Line<F>>, ..) 改成 fn gelu_array<F: Float, N: Size>(input: &Array<Vector<F, N>>, ..)，启动参数从 ArrayArg::from_raw_parts::<f32>(input_handle, input.len(), line_size) 变成在 launch 上单独传 vector_size：宽度成了类型的一部分，进 KernelId，同一个宽度编一次。#1322（「mega-refactor: Totally change the frontend to enable references, among other things」，首发 v0.11.0-pre.1）再把参数换成 input: &[Vector<F, N>]，ArrayArg 换成 BufferArg，同一个 diff 里 x * (Vector::erf(tmp) + 1.0) / 2.0 也改成了显式的 Vector::one() 和 Vector::new(F::new(2.0f32))。两个提交只有标题、没有正文。所以 0.22.1 的 burn 拉进来的 cubecl，和 2026 年初以前的 cubecl 教程，内核签名写法互不相通。

launch：在调用线程上只做登记

展开出的 saxpy::launch（生成它的代码）原样是：

    pub fn launch<'kernel, F: Float>(
        __client: &cubecl::prelude::Client,
        __cube_count: cubecl::prelude::CubeCount,
        __cube_dim: cubecl::prelude::CubeDim,
        _n: usize,
        a: cubecl::RuntimeArg<f32>,
        x: cubecl::RuntimeArg<[Vector<F, N>]>,
        y: cubecl::RuntimeArg<[Vector<F, N>]>,
    ) {
        let mut __settings = cubecl::prelude::KernelSettings::new(
            __cube_dim.into(),
            cubecl::prelude::ExecutionMode::Checked,
            cubecl::prelude::AddressType::U32,
        );
        let mut launcher = cubecl::prelude::KernelLauncher::new(__settings.clone());
        launcher
            .with_scope(|scope| {
                scope.device_properties(__client.properties());
                scope
                    .register_type::<
                        cubecl::prelude::DynamicScalar<_F>,
                    >(F::elem_type(&scope));
                scope.register_size::<cubecl::prelude::DynamicSize<_N>>(_n.into());
            });
        launcher.declare_io(cubecl::prelude::BufferIOAttr::ReadWrite);
        let comp_arg_0 = <f32 as cubecl::prelude::LaunchArg>::register(a, &mut launcher);
        launcher.declare_io(cubecl::prelude::BufferIOAttr::ReadOnly);
        let comp_arg_1 = <[Vector<
            F,
            N,
        >] as cubecl::prelude::LaunchArg>::register(x, &mut launcher);
        launcher.declare_io(cubecl::prelude::BufferIOAttr::ReadWrite);
        let comp_arg_2 = <[Vector<
            F,
            N,
        >] as cubecl::prelude::LaunchArg>::register(y, &mut launcher);
        let __kernel = Saxpy::<
            F,
        >::new(
            __settings,
            __client.properties_shared(),
            __client.target_properties_shared(),
            comp_arg_0,
            comp_arg_1,
            comp_arg_2,
            _n,
        );
        launcher.launch(__cube_count, __kernel, __client)
    }

每个参数的 register 把运行期的部分（句柄的绑定、标量的字节、缓冲长度）记进 KernelLauncher，返回编译期的部分（CompilationArg：缓冲是否原地复用某个输入、标量没有）。宏从参数的可变性推出读写属性：x 是 &[..]，标成 ReadOnly；y 是 &mut [..]，标成 ReadWrite。Saxpy::new 只收编译参数，不碰数据。KernelLauncher::launch 把绑定打包、把 Saxpy 装进 Box<dyn CubeKernel>，交给 client.launch。这个调用把任务丢给设备线程就返回（见运行时与 Client），整个 launch 没有构造一条 IR。

IR 在设备线程上构造。wgpu 运行时的 WgpuServer::pipeline 开头：

    fn pipeline(
        &mut self,
        kernel: Box<dyn CubeKernel>,
        bindings: &KernelArguments,
    ) -> Result<PipelineEntry, LaunchError> {
        let kernel_id = kernel.id();
        let mode = kernel_id.mode;

        if let Some(pipeline) = self.pipelines.get(&kernel_id) {
            return Ok(pipeline.clone());
        }
        // ...
        let definition = cubecl_core::define_kernel(&*kernel)?;
        // ...
        let mut compiler = C::init(self.backend, &self.compilation_options);
        let mut compiled = compiler.compile_kernel(self, kernel, definition)?;

先用 kernel.id() 查已经建好的管线，命中就直接用；没命中才 define_kernel → Saxpy::define → expand，再交给编译器。define_kernel 用 catch_unwind 包住 define，expand 里的 panic（比如两个宽度不同的向量相加）变成 CompilationError::ExpansionPanicked，挂在这次启动要写的缓冲上，读那个缓冲时才报出来。编译和管线缓存的其余部分见 wgpu 与 Metal。

看生成的 MSL

cubecl 的调试日志会把每个编译出的内核源码打出来：环境变量 CUBECL_DEBUG_LOG=stdout（由 override_from_env 读取）把编译、性能、autotune 三种日志都打到标准输出。跑 saxpy，98 行输出，内核部分（typedef 几行省略，标 ...；日志在 MSL 源码前后各有一行 Markdown 代码围栏，这里删掉了）：

[START_KERNEL_COMPILATION]
name: cube_kernels::saxpy::Saxpy<f32>
id: KernelId {
    type: cube_kernels::saxpy::Saxpy<f32>,
    address_type: U32,
    cube_dim: Dim3 {
        x: 4,
        y: 1,
        z: 1,
    },
    mode: Checked,
    info: SaxpyInfo {
        n: 4,
        a: (),
        x: BufferCompilationArg {
            inplace: None,
        },
        y: BufferCompilationArg {
            inplace: None,
        },
    },
}
source:

#include <metal_stdlib>
using namespace metal;
...
struct alignas(16) float_4 { float i_0; float i_1; float i_2; float i_3; };

struct info_st {
    float scalars_f32[2];uint32_t static_meta[2];
};
[[max_total_threads_per_threadgroup(64)]] [[kernel]] void saxpy_f_f32_n_4_daf77f4d(device float_4 const* const v2 [[buffer(0)]], device float_4 * const v6 [[buffer(1)]], constant info_st const* const v33, uint3 const threadgroups_per_grid_v166 [[threadgroups_per_grid]], uint32_t const thread_index_in_threadgroup_v167 [[thread_index_in_threadgroup]], uint3 const threadgroup_position_in_grid_v168 [[threadgroup_position_in_grid]]) {
constant uint32_t const* const dynamic_meta_v35 = reinterpret_cast<constant uint32_t const*>(v33 + 1);
constant info_st& info_v34 = *v33;
float const v36 = info_v34.scalars_f32[0];
uint32_t const v37 = info_v34.static_meta[0];
uint32_t const v169 = threadgroup_position_in_grid_v168[2];
...
uint32_t const v162 = 4;
uint32_t const v85 = v82 * v162;
uint32_t const v87 = uint32_t(thread_index_in_threadgroup_v167);
uint32_t const v88 = v85 + v87;
bool const v12 = v88 < v37;
if(v12) {
device float_4 const* const v18 = &v2[v88];
float_4 const v19 = *v18;
float const v41 = v19.i_0;
float const v42 = v36 * v41;
...
float const v63 = v26.i_3;
float const v64 = v51 + v63;
float_4 const v65 = float_4{v55, v58, v61, v64};
*v25 = v65;
}
return;
}

[END_KERNEL_COMPILATION]

| 12.334µs   | Saxpy | f32_daf77f4d

对着 Rust 源码读：

  • KernelId 就是缓存键，info: SaxpyInfo { n: 4, .. } 里的 n 是向量宽度。a: () 说明标量没有编译期部分：换一个 a 不会重新编译。
  • 入口名 saxpy_f_f32_n_4_daf77f4d 由 Saxpy::new 拼出：函数名、_f_f32（F 的类型名）、_n_4（宽度），再加一个哈希。
  • ABSOLUTE_POS 没有对应的 Metal 内建，由 threadgroup_position_in_grid、threadgroups_per_grid 和 thread_index_in_threadgroup 算出来（v88）。乘的 4 是 cube_dim.x，它是编译期常量（KernelId 里有 cube_dim）。
  • x.len() 是 info_st.static_meta[0]（v37）；标量 a 是 info_st.scalars_f32[0]（v36）。启动时的元数据和标量都放进同一个常量缓冲 info_st。
  • Vector<f32, 4> 编成了 float_4 结构体，乘加逐个分量写出（v19.i_0 到 i_3），没有用 Metal 原生的 float4。
  • 我们写的 if ABSOLUTE_POS < x.len() 就是 if(v12)。

最后一行 | 12.334µs | Saxpy | 是性能日志：CUBECL_DEBUG_LOG 同时把性能日志调到 Medium，Client::launch_inner 于是走「先测时间再提交」的分支，每次启动都要等设备做完（submit_blocking），不再是丢给设备线程就返回。开着它测出的时间不能当正常运行的时间。

同一个函数，几个内核

KernelId 由内核类型、地址类型、cube_dim、执行模式和 Info 组成（struct KernelId）。第二个实验内核把这几样各变一下：

/// out = x + y with no bounds guard of our own; `relu` is decided when the kernel is compiled.
#[cube(launch, launch_unchecked)]
fn add<F: Float>(x: &[F], y: &[F], out: &mut [F], #[comptime] relu: bool) {
  let v = x[ABSOLUTE_POS] + y[ABSOLUTE_POS];
  if relu {
    out[ABSOLUTE_POS] = F::max(v, F::new(0.0f32));
  } else {
    out[ABSOLUTE_POS] = v;
  }
}

relu 带 #[comptime]，展开后 expand 的签名里它就是一个普通的 relu: bool，函数体里那个 if 也原样留成 Rust 的 if，没有变成 if_expand：

            if relu {
                {
                    let _value = cubecl::prelude::IntoExpand::into_expand(
                        {
                            F::__expand_max(
                                scope,
                                (v).into(),
                                ({ F::__expand_new(scope, (0.0f32).into()) }).into(),
                            )
                        },
                        scope,
                    );
                    // ...

所以 relu = true 时 IR 里只有 max 那一支，false 时只有赋值那一支；relu 的值同时进了 AddInfo，两个值是两个 KernelId。

❗IMPORTANT
动脑筋：同一个 add，启动四次，编译几次？

run_add(client, checked, relu) 用 6 个元素、8 个单元调一次 add，checked 选 launch 还是 launch_unchecked。程序依次调用 (true, false)、(false, false)、(true, true)、(true, false)，开 CUBECL_DEBUG_LOG=stdout，数 [START_KERNEL_COMPILATION] 出现几次。答案在本页末尾。

launch 与 launch_unchecked

#[cube(launch, launch_unchecked)] 生成两个函数，区别只在交给 KernelSettings 的执行模式：launch 是 ExecutionMode::Checked，launch_unchecked 是 unsafe fn、ExecutionMode::Unchecked。按设计，Checked 模式下编译器管线里的 CheckedIo pass 把每个带 checked 标记的下标改写成 checked_index：

pub fn checked_index(index: usize, buffer_len: usize) -> usize {
    index.min(buffer_len - 1)
}

切片的 x[i] 确实带这个标记（IndexExpand for SliceExpand 调 read_offset::expand(.., true)），MSL 编译器也确实挂了这个 pass（cubecl-cpp 的 add_pass(CheckedIoPass::new(..))）。但上面 add 的 MSL 里，launch 和 launch_unchecked 编出来的两份除了入口名的哈希和内建参数的顺序，逐行一样，都没有 min：

device float const* const v17 = &v1[v60];
float const v18 = *v17;
device float const* const v23 = &v5[v60];
float const v24 = *v23;
float const v25 = v18 + v24;
device float * const v30 = &v9[v60];
*v30 = v25;

再跑一个越界读的程序：x、y 各 6 个元素，out 8 个，启动 8 个单元，单元 6、7 读 x[6]、x[7]。如果下标被夹到 len - 1，它们应当读到 x[5] + y[5] = 66：

oob checked=true out = [11.0, 22.0, 33.0, 44.0, 55.0, 66.0, 0.0, 0.0]
oob checked=false out = [11.0, 22.0, 33.0, 44.0, 55.0, 66.0, 0.0, 0.0]

两种模式都读到了 0（分配的对齐填充），没有夹取。原因在 pass 的匹配条件。CheckedIo 只改写 base 本身是 RuntimeArrayType 的 IndexOp：

impl MatchRewrite for CheckedIo {
    fn r#match(&mut self, ctx: &Context, op: Ptr<Operation>) -> bool {
        Operation::get_op::<IndexOp>(op, ctx)
            .is_some_and(|it| it.checked(ctx) && is_runtime_array(ctx, it.base(ctx)))
    }
    // ...
}

fn is_runtime_array(ctx: &Context, value: Value) -> bool {
    let ty = value.get_type(ctx).deref(ctx);
    ty.downcast_ref::<RuntimeArrayType>().is_some()
}

而全局缓冲在 Scope::global 里的类型是 PointerType::get(ctx, ty_arr.into(), AddressSpace::Global(..))，即「指向」运行期数组的指针，IndexOp 的 base 也要求是指针（indexed_ptr_ty 里 .expect("Should be pointer")）。外层是 PointerType，downcast_ref::<RuntimeArrayType>() 落空，这个 pass 对全局缓冲不会触发。这是本卷在 0.11.3 + MSL 上实测并读源码得出的结论，CUDA / SPIR-V 编译器挂的是同一个 pass，只读了源码、没有实测。实际影响是：在这个版本里不能指望 launch 把越界下标夹回来，内核要自己写 if ABSOLUTE_POS < x.len() 这样的守卫。

📝NOTE
冷知识：越界检查原来是一个指令访问器

#1402（「refactor: Make everything Pliron」，首发 v0.11.0-pre.3）把 cubecl 自己的 IR 整个换成了 pliron（一个用 Rust 写的、MLIR 风格的 IR 框架），提交只有标题。改之前，index_expand 自己建输出值、登记一条 Instruction：scope.register(Instruction::new(Memory::Index(IndexOperands { list, index, unroll_factor: 1, checked }), output))，输出类型写成 Type::pointer(ty, class)；越界检查是一个 CheckedIoVisitor，在 Memory::Index(op) if op.checked 且 op.list.has_buffer_length() 时展开成带长度的版本。改之后，同一个函数变成上面两行的 IndexOp::maybe_checked + register_with_result，检查变成 pliron 的 MatchRewrite，条件换成了按 base 的类型判断，也就是上一节对全局缓冲不触发的那个条件。

动脑筋答案

3 次。 日志里 [START_KERNEL_COMPILATION] 出现 3 次，第四次调用没有编译、直接复用了第一次的管线（入口名的哈希也一样）：

3
name: cube_kernels::add::Add<f32>
    mode: Checked,
        relu: false,
[[max_total_threads_per_threadgroup(64)]] [[kernel]] void add_f_f32_7c592e50
| 12.25µs    | Add | f32_7c592e50
add checked=true relu=false out = [1.5, -1.5, 3.5, -3.5, 5.5, -5.5]
name: cube_kernels::add::Add<f32>
    mode: Unchecked,
        relu: false,
[[max_total_threads_per_threadgroup(64)]] [[kernel]] void add_f_f32_b1476a3d
| 12.667µs   | Add | f32_b1476a3d
add checked=false relu=false out = [1.5, -1.5, 3.5, -3.5, 5.5, -5.5]
name: cube_kernels::add::Add<f32>
    mode: Checked,
        relu: true,
[[max_total_threads_per_threadgroup(64)]] [[kernel]] void add_f_f32_a1caed20
| 13.166µs   | Add | f32_a1caed20
add checked=true relu=true out = [1.5, 0.0, 3.5, 0.0, 5.5, 0.0]
add checked=true relu=false out = [1.5, -1.5, 3.5, -3.5, 5.5, -5.5]
| 12.167µs   | Add | f32_7c592e50

（第一行是 grep -c 的计数，其余是 grep 出的入口名、模式、relu 和结果行。）

  • 第二次只换了执行模式：mode 是 KernelId 的字段，Unchecked 是另一个键，所以重新编译，尽管上一节看到它编出的 MSL 和 Checked 一样。
  • 第三次只换了 relu：#[comptime] 参数进了 AddInfo，也是另一个键，编出的 MSL 多了 max(v25, v26)。
  • 第四次和第一次的 KernelId 完全相同，WgpuServer::pipeline 在 self.pipelines 里命中，连 define 都没调。输入数据不同不影响：数据只在运行期参数里，不在键里。

小结

  • #[cube(launch)] 保留原函数只做类型检查，另生成同名模块：expand（构造 IR）、内核描述体 Saxpy<F>（只带编译参数）、KernelMetadata::id、CubeKernel::define、launch。
  • expand 把每个运算换成 __expand_* 调用，最后都是建一个 pliron 操作、Scope::register 插进去；泛型 F、N 换成 DynamicScalar / DynamicSize，真实类型查 scope 的类型表。
  • 0.11 的前端：全局缓冲是 &[T] / &Tensor<T>，向量是 Vector<F, N>，Array 只剩内核里的定长数组；#[comptime] 参数在展开时就是普通 Rust 值，并进入 KernelId。
  • launch 在调用线程上只登记参数、构造描述体、提交；IR 在设备线程上、按 KernelId 第一次编译时才构造，之后复用管线。
  • CUBECL_DEBUG_LOG=stdout 能看到每个内核的 MSL，但它同时让每次启动变成阻塞的，不能用来计时。
  • 0.11.3 的 MSL 路径上，launch 和 launch_unchecked 编出相同的代码，越界不会被夹取；内核要自己写下标守卫。

💬 关联讨论

发起讨论
加载中...