Rust in Everything
基于 burn 0.22.1

wgpu 与 Metal

主线位置：算子主线最后一段「cubecl 编译 / 提交」在 macOS 上的落地。CubeBackend 和 cubek 把内核交给 Client::launch 之后（见运行时与 ComputeClient），wgpu 运行时的服务端 WgpuServer 要做三件事：第一次用到设备时选适配器、决定用哪种着色语言；第一次遇到某个内核时把 cubecl IR 编成 MSL、交给 Metal 建管线；之后每次提交直接用缓存的管线。本页讲这三件事。内核本身怎么写成 IR 见 #[cube] 内核。

flowchart TB
  dev["Device::metal(kind)"] --> wd["WgpuDevice backend Metal"]
  wd --> setup["try_create_setup 选适配器"]
  setup --> feat["register_metal_features"]
  feat -->|"MSL 3.2 可用"| msl["AutoCompiler::Msl"]
  feat -->|不可用| wgsl["AutoCompiler::Wgsl"]
  launch["Client::launch 内核"] --> pipe["WgpuServer::pipeline"]
  pipe -->|命中| cached["进程内管线表"]
  pipe -->|未命中| ir["define_kernel 得 IR"]
  ir --> src["MslCompiler 生成 MSL 源码"]
  src --> mod["create_shader_module_passthrough"]
  mod --> mtl["Metal newLibraryWithSource"]
  mtl --> cpl["create_compute_pipeline"]
  cpl --> cached

本页源码是 cubecl 0.11.3（cubecl-wgpu、cubecl-server、cubecl-runtime）和 burn 0.22.1。实验 crate 有三个，都在 exp-burn/ 下：rt-tune（burn 只开 metal、tracing）、rt-wgsl（只开 wgpu、tracing）、rt-cube（直接依赖 burn-cubecl）。平台 M2 Max、macOS 27、rustc 1.99.0，计时都是 release 构建。Metal 是实测；Vulkan（SPIR-V）和 WebGPU（WGSL 在浏览器里）只读了源码。

metal feature 走的是哪条路

burn 的 metal 先打开 wgpu，再一路传到 burn-wgpu 的 metal = ["cubecl/metal", ..]。到了 cubecl 这一层：

metal = ["wgpu-msl"]
metal-native = ["cubecl-metal"]
# ...
wgpu-msl = ["wgpu", "cubecl-wgpu/msl"]

metal = ["wgpu-msl"]：还是 wgpu 运行时，只是给它的 msl feature 打开了 MSL 编译器（cubecl-cpp 的 metal 方言）。cubecl 另有一个原生 Metal 运行时 cubecl-metal，feature 名是 metal-native，burn 没有任何 feature 通向它。rt-tune 的依赖树可以对一下：

cubecl-wgpu feature "cubecl-cpp"
cubecl-wgpu feature "default"
cubecl-wgpu feature "msl"
cubecl-wgpu feature "sanitize-filename"
cubecl-wgpu feature "std"
cubecl-wgpu feature "tracing"
cubecl-metal / cubecl-spirv in tree: 0
cubecl-cpp in tree: 1

（cargo tree -e features -i cubecl-wgpu，去掉了重复行和树线。）所以一次 matmul 里打印的后端名是 cubecl<wgpu<msl>>，名字里的 msl 来自 runtime_name：

fn runtime_name(backend: wgpu::Backend, options: &WgpuCompilationOptions) -> &'static str {
    match backend {
        wgpu::Backend::Vulkan if options.supports_vulkan_compiler => "wgpu<spirv>",
        wgpu::Backend::Metal if options.supports_msl_compiler => "wgpu<msl>",
        _ => "wgpu<wgsl>",
    }
}

它看的是运行期的 supports_msl_compiler，不是编译期的 feature：开了 metal、但设备撑不起 MSL 时，名字会是 wgpu<wgsl>（下一节）。

📝NOTE
没有蠢问题

问： cubecl 自己也有 Device::metal，它和 burn 的 Device::metal 是一回事吗？

答： 不是，名字撞了。cubecl 的 Device 枚举里 Metal(MetalDevice) 是原生运行时 cubecl-metal 的设备，Wgpu(WgpuDevice) 才是 wgpu 的。burn 的 Device::metal(kind) 对应的是 cubecl 的 Device::metal_msl，它的文档注释写着「Not the native Metal runtime, which is [Device::metal] and a different device altogether」。直接写 cubecl 代码时用错了，client() 会 panic 在「belongs to a runtime this build does not link」。

选设备：DeviceKind 和图形 API

burn 的四个 wgpu 构造函数共用 wgpu_device，把 DeviceKind 映射成 cubecl 的 WgpuDeviceKind，再用 .on(backend) 定图形 API（映射本身见 Device）。Device::wgpu 给 WgpuBackend::Auto，Device::metal 给 WgpuBackend::Metal。构造是惰性的：真正去找适配器是第一次用到这个设备时，在 cubecl 的设备服务线程上调 DeviceService::init → try_create_setup。

try_create_setup 先按 backend_candidates 列出要试的图形 API：固定了 API 的只有它自己一个；Auto 交给 AutoGraphicsApi::chain：

        cfg_if::cfg_if! {
            if #[cfg(target_family = "wasm")] {
                alloc::vec![Backend::BrowserWebGpu]
            } else if #[cfg(target_os = "macos")] {
                alloc::vec![Backend::Metal]
            } else {
                alloc::vec![Backend::Vulkan, Backend::Dx12, Backend::Gl]
            }
        }

macOS 上 Auto 只有 Metal 一个候选，所以 Device::wgpu(..) 和 Device::metal(..) 最后用的是同一个 Metal 适配器。挑适配器看 kind（request_adapter）：

DeviceKind怎么挑
DefaultDevicerequest_adapter 带 PowerPreference::HighPerformance；环境变量 CUBECL_WGPU_DEFAULT_DEVICE 可改成别的种类（try_device_override，写法如 IntegratedGpu(1)、Cpu）
DiscreteGpu(i) / IntegratedGpu(i) / VirtualGpu(i) / Cpu枚举这个 API 的全部适配器，只留 wgpu 报告的种类相同的，取第 i 个；没有就报 No Discrete GPU device found 之类，后面附上找到的适配器列表
Existing(id)不找适配器，指向用 Device::wgpu_options().setup(..) 登记过的外部 wgpu 设备

M2 Max 上只有一个 Metal 适配器，wgpu 报它是 IntegratedGpu。rt-tune kind <种类> 在 Device::metal(kind) 上做一次加法（release）：

$ rt-tune kind integrated
Device<Cube(Wgpu(WgpuDevice { kind: IntegratedGpu(0), backend: Metal }))>
[2.0, 4.0] name=dispatch<fusion<cubecl<wgpu<msl>>>>
$ rt-tune kind discrete
Device<Cube(Wgpu(WgpuDevice { kind: DiscreteGpu(0), backend: Metal }))>

thread 'DSD-0-16384' (3354333) panicked at ~/.cargo/registry/src/index.crates.io-1949cf8c6b5b557f/cubecl-wgpu-0.11.3/src/runtime.rs:51:57:
Unable to acquire wgpu device: AdapterUnavailable { device: WgpuDevice { kind: DiscreteGpu(0), backend: Metal }, message: "No Discrete GPU device found, adapters [AdapterInfo { name: \"Apple M2 Max\", vendor: 0, device: 0, device_type: IntegratedGpu, device_pci_bus_id: \"\", driver: \"\", driver_info: \"\", backend: Metal, subgroup_min_size: 4, subgroup_max_size: 64, transient_saves_memory: Some(true), limit_bucket: None }]" }
note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace

thread 'main' (3354329) panicked at ~/.cargo/registry/src/index.crates.io-1949cf8c6b5b557f/cubecl-common-0.11.3/src/device/handle/channel.rs:409:35:
called `Result::unwrap()` on an `Err` value: RecvError

IntegratedGpu(1) 和 Cpu 也是同样的两段 panic。适配器是在设备服务线程（DSD-<type_id>-<index_id>）上找的，找不到时 DeviceService::init 的 expect("Unable to acquire wgpu device") 让那个线程 panic，有用的信息在这一段；main 线程随后只看到通道断开的 RecvError。在 Apple 芯片上写 DiscreteGpu(0) 就会这样，用 DefaultDevice 或 IntegratedGpu(0)；想拿到 Result 而不是 panic，用 Device::wgpu_options().device_kind(..).init()。

两条路落到同一块 GPU，差别在别处。rt-tune devices 依次在 Device::default()、Device::wgpu(DefaultDevice)、Device::metal(DefaultDevice) 上做一次加法（autotune 缓存已有，加法没有调优）：

Device::default()              Device<Cube(Wgpu(WgpuDevice { kind: DefaultDevice, backend: Auto }))>
                               name=dispatch<fusion<cubecl<wgpu<msl>>>> first op 195.649042ms -> [2.0, 4.0]
Device::wgpu(DefaultDevice)    Device<Cube(Wgpu(WgpuDevice { kind: DefaultDevice, backend: Auto }))>
                               name=dispatch<fusion<cubecl<wgpu<msl>>>> first op 9.282583ms -> [2.0, 4.0]
Device::metal(DefaultDevice)   Device<Cube(Wgpu(WgpuDevice { kind: DefaultDevice, backend: Metal }))>
                               name=dispatch<fusion<cubecl<wgpu<msl>>>> first op 49.621625ms -> [2.0, 4.0]
wgpu == metal: false
  • 三个都编成了 MSL（wgpu<msl>）。
  • Device::wgpu 和 Device::default() 是同一个设备（开了 metal 时默认设备就是 Auto，见 Device），第二次只用 9 ms。Device::metal 又花了 50 ms：它的 DeviceId 不同（图形 API 编在 id 里），cubecl 为它另起了一套服务线程、内存池和管线表，同一个加法内核要再编译一次。
  • 还有两处不同。一是 MSL 不可用时，Auto 允许退回 WGSL，固定成 Metal 的不允许，初始化时 panic（下一节）。二是 autotune 的结果按设备分开存，换一种写法就要重新调优，见本页的动脑筋。

MSL 能不能用：先编一个金丝雀

设备拿到之后，try_create_server 调 register_features：先问 Vulkan（只有开了 spirv 才可能是），再问 Metal，都不是就按 WGSL 登记设备能力。Metal 这一步是 register_metal_features：

    // Canary for ensuring the current Metal language version is 3.2+.
    const LAMBDA_CANARY: &str = r#"
        #include <metal_stdlib>
        using namespace metal;

        kernel void __canary__(device float* out [[buffer(0)]]) {
            auto f = [](float x) -> float { return x * 2.0; };
            out[0] = f(1.0);
        }
    "#;
    // ...
        let has_required_family = raw.supportsFamily(MTLGPUFamily::Metal3)
            || raw.supportsFamily(MTLGPUFamily::Apple7)
            || raw.supportsFamily(MTLGPUFamily::Mac2);

        let support = MslSupport::new(has_required_family, || {
            autoreleasepool(|_| {
                let canary = NSString::from_str(LAMBDA_CANARY);
                let options = MTLCompileOptions::new();
                options.setLanguageVersion(MTLLanguageVersion::Version3_2);
                raw.newLibraryWithSource_options_error(&canary, Some(&options))
                    .map(|_| ())
                    .map_err(|err| format!("{err:?}"))
            })
        });

两道检查：GPU 家族要有 SIMD-scoped 运算（plane 和 CMMA 内核靠它），系统的 Metal 编译器要接受 MSL 3.2（生成的源码里用了 lambda）。都过了才把 supports_msl_compiler 设成 true，再登记类型（f16、bf16、i64 都可算，fp8 / fp4 只能存和转换；原子只登记了 i32、u32 和 f32 的加法与读写，注释说 64 位原子登记过、结果算错，又撤掉了）和 CMMA 尺寸。任一道不过，只打一条 log::warn!，设备按 WGSL 继续。

接着 AutoCompiler::validate_runtime 决定能不能接受这种退回：只有请求的是 WgpuBackend::Metal 时才要求 MslCompiler 可用，否则报 WgpuCompilerInitError::Unavailable（「use AutoCompiler with an automatic device to allow WGSL fallback」）。走惰性初始化时这个错误被 expect("Unable to initialize wgpu runtime") 变成 panic；用 Device::wgpu_options().graphics_api(WgpuBackend::Metal).init() 则拿到 Err。M2 Max 两道都过，本页的实验看不到退回。

📝NOTE
冷知识：MSL 曾经是无条件打开的

PR #1512（「Add version check to ensure required features are supported in MSL」，cubecl 首次进入 v0.11.0-pre.3）之前，register_metal_features 只要拿到 Metal 适配器就调 register_features，没有返回值，也不检查什么：开了 msl 的构建在任何 Metal 设备上都生成 MSL。生成的源码用了 MSL 3.2 才有的 lambda，系统编译器旧一点就编不过。这个 PR 加了 lambda 金丝雀和 Apple7 家族检查，函数改成返回 bool，不支持时退回 WGSL 并 log::warn!（提交说明里的「Actually disable the compiler if unsupported」）。随后的 PR #1540（「fix(metal): correct MSL capability detection」，同样首次进入 v0.11.0-pre.3）把家族条件放宽成上面的 Metal3 || Apple7 || Mac2，并给金丝雀显式指定 Version3_2；现在的注释说明这和 wgpu-hal 判断 SIMD-scoped 能力的条件一致。

编一个内核：从 IR 到管线

每次提交内核，WgpuServer 先拿 pipeline 找管线：

        let kernel_id = kernel.id();
        let mode = kernel_id.mode;

        if let Some(pipeline) = self.pipelines.get(&kernel_id) {
            return Ok(pipeline.clone());
        }

        let mut recording = CompilationRecording::new(&kernel_id);
        let cached = self.load_cached_pipeline(&kernel_id, bindings, mode)?;
        // ...
        let definition = cubecl_core::define_kernel(&*kernel)?;
        recording.defined(&definition);

        let mut compiler = C::init(self.backend, &self.compilation_options);
        let mut compiled = compiler.compile_kernel(self, kernel, definition)?;
        // ...
        let module = self.create_module(
            &compiled.entrypoint_name,
            kernel_id.cube_dim.into(),
            ModuleSource::resolve(repr, compiler.lang_tag(), &compiled.source)?,
            mode,
        )?;
        let pipeline = self.create_pipeline(&compiled.entrypoint_name, repr, module, bindings)?;
        self.pipelines.insert(
            kernel_id.clone(),
            (pipeline.clone(), compiler_info, io.clone()),
        );
  1. 查表。 键是 KernelId：内核类型加上它的编译期参数（cube_dim、向量宽度、dtype、是否原地写等）。同一个内核换了编译期参数就是另一条，要重新编。
  2. 展开成 IR。 define_kernel 运行 #[cube] 宏生成的展开代码，得到 cubecl IR（#[cube] 内核）。
  3. IR 到源码。 AutoCompiler::init 按上一节的 supports_msl_compiler 选 Msl，MslCompiler（cubecl-cpp 的 Metal 方言）把 IR 写成 MSL 文本。
  4. 源码到 Metal 库。 create_module 对 MSL 走 create_shader_module_passthrough，把文本原样交给 wgpu；wgpu-hal 的 Metal 后端（wgpu-hal-30.0.1 的 src/metal/device.rs）收到 ShaderInput::Msl 就调 newLibraryWithSource_options_error，由系统的 Metal 编译器当场编译。整个过程没有 naga 参与，wgpu 也不校验源码。
  5. 建管线。 create_pipeline 按 metal::bindings 建绑定布局：所有缓冲都标成读写（同一块显存可能被多个切片共享，分不清谁只读），元数据缓冲是 uniform 或只读，然后 create_compute_pipeline。模块和管线的创建都包在 validated 里：wgpu 的验证错误和内部错误被收成这次启动的 CompilationError，不会变成轮询线程上的 panic。

CUBECL_DEBUG_LOG=stdout 会把第 3 步的结果打出来。rt-tune add 在 Metal 上把两个长 4 的张量相加，输出里只有一个内核（截掉了中间的下标计算）：

[START_KERNEL_COMPILATION]
name: burn_cubecl::kernel::binary::kernel_binop::KernelBinop<
    burn_cubecl::kernel::binary::AddOp,
>
id: KernelId {
    type: burn_cubecl::kernel::binary::kernel_binop::KernelBinop<burn_cubecl::kernel::binary::AddOp>,
    address_type: U32,
    cube_dim: Dim3 {
        x: 32,
...
        out: ArrayView {
            buffer: BufferCompilationArg {
                inplace: Some(
                    0,
                ),
...
source:
```msl

#include <metal_stdlib>
using namespace metal;
...
struct alignas(16) float_4 { float i_0; float i_1; float i_2; float i_3; };
...
[[max_total_threads_per_threadgroup(64)]] [[kernel]] void kernel_binop_c_f32_n_4_2f077c2a(device float_4 * const v2 [[buffer(0)]], device float_4 const* const v7 [[buffer(1)]], constant info_st const* const v41, ...) {
...
if(v22) {
device float_4 * const v28 = &v2[v86];
float_4 const v29 = *v28;
device float_4 const* const v34 = &v7[v86];
float_4 const v35 = *v34;
float const v51 = v29.i_0;
float const v52 = v35.i_0;
float const v53 = v51 + v52;
...
float_4 const v63 = float_4{v53, v56, v59, v62};
*v28 = v63;
}
return;
}

几处可以和源码对上：KernelId 里的 inplace: Some(0) 是 launch_binop 发现左操作数可以原地写（burn-cubecl 后端），所以 MSL 里只有两个缓冲、结果写回 v2；n_4 和 float_4 是向量宽度 4；入口名末尾的 2f077c2a 是编译期参数的哈希，换参数就换名字。

对照：同一块 GPU 上的 WGSL

只开 wgpu、不开 metal 时，cubecl-wgpu 不带 msl feature，macOS 上 Auto 仍然是 Metal API，但编译器只能是 WgslCompiler。rt-wgsl 用 Device::wgpu(DefaultDevice) 跑同一个加法，日志里的源码是 WGSL：

source:
```wgsl
@group(0) @binding(1) var<storage, read_write> v7: array<vec4<f32>>;
@group(0) @binding(0) var<storage, read_write> v2: array<vec4<f32>>;
struct info_st { scalars_usize: array<u32, 4>, static_meta: array<u32, 2> }
@group(0) @binding(2) var<storage, read> info_var: info_st;
@compute @workgroup_size(32, 1, 1)
fn kernel_binop_c_f32_n_4_91757fd7(@builtin(workgroup_id) v137: vec3<u32>, @builtin(num_workgroups) v139: vec3<u32>, @builtin(local_invocation_index) v146: u32) {
...

这条路多两道工序：create_module 走 create_shader_module_trusted（关掉 wgpu 的越界检查，cubecl 自己负责），wgpu-hal 在建管线时用 naga 把 WGSL 翻成 MSL，再交给同一个 newLibraryWithSource。翻译之外更大的差别在能力：WGSL 登记不了 CMMA，matmul 的 autotune 只能在不用矩阵单元的候选里挑。两边各跑一次空缓存的进程，看 autotune 记下的结果（[autotune.decisions] stdout = true，见 autotune）和之后 20 次的最短时间：

后端名1024² 胜出的候选（中位数）2048² 胜出的候选（中位数）2048² 之后 20 次最短测出的计算峰值
wgpu<msl>matmul_simple_cyclic_cmma_multi_rows（1.33 ms）同左（3.12 ms）3.6–3.8 ms9.04 × 10¹² op/s（ComputeCmma）
wgpu<wgsl>matmul_double_unit_min_tile_size（1.52 ms）同左（13.86 ms）13.5–14.3 ms2.76 × 10¹² op/s（ComputeDirect）

1024² 两边差不多，2048² 上 MSL 快三倍多。最后一列是 autotune 自己测的峰值（见 autotune），它说明了差距从哪来：MSL 路径用得上 SIMD-group 矩阵运算。数字只对这台机器有效。

编译缓存：只在进程里

上面的管线表 pipelines 是 WgpuServer 的一个字段，进程结束就没了。WgpuServer::new 里决定它背后有没有持久化：

        #[cfg(feature = "spirv")]
        let spirv_cache = compilation_store(
            "vulkan",
            format!("spirv_{}_{}", adapter_info.vendor, adapter_info.device),
        );

        // WGSL is compiled by the driver on every run, so without the SPIR-V
        // store there is nothing persisted for a switch to invalidate.
        #[cfg(feature = "spirv")]
        let pipelines = CompilationCache::mirroring(&spirv_cache);
        #[cfg(not(feature = "spirv"))]
        let pipelines = CompilationCache::unbound();

持久的编译缓存只给 SPIR-V 准备：pipeline 末尾的注释是 「Only a SPIR-V kernel is stored: any other build changes nothing.」。compilation_store 还要再看配置项 compilation.cache（默认 false），两者都满足才在 target/environment/default.db 里开一个 vulkan 命名空间。原生的 CUDA、HIP、cubecl-metal 运行时各有自己的编译缓存（命名空间 cuda、hip、metal）；wgpu<msl> 没有。建管线时传给 wgpu 的 cache: None 也说明没有用 wgpu 的 PipelineCache。

一次 matmul 引了 burn 文档「Compiled-kernel caching also requires compilation.cache = true」。在 rt-tune 目录下放一个 cubecl.toml：

[compilation]
cache = true

先确认配置读到了（rt-tune config 打印 CubeClRuntimeConfig::get()），再连跑两个进程，各做一次 1024² matmul（autotune 缓存已有）：

compilation.cache = false, autotune.level = Balanced, autotune.disable_cache = false
compilation.cache = true, autotune.level = Balanced, autotune.disable_cache = false
--- 进程 1
1024x1024 first:    133.527ms  launches=1 [MatmulEntry=1, ~compile(integrate)=1]
--- 进程 2
1024x1024 first:    687.952ms  launches=1 [MatmulEntry=1, ~compile(integrate)=1]

（~compile(integrate) 数的是 cubecl 的 integrate span，每编一个新内核一次。）第二个进程仍然编译了一次，数据库里也没有多出编译相关的命名空间（sqlite3 列出的五个命名空间和打开前一样：一个 autotune/..、一个 throughput/..、三个 records/v1/..）。在 wgpu<msl> 上打开 compilation.cache 什么也不改变，每个新进程的第一次 matmul 都要重新编译 MSL。两次第一次调用 133 ms 和 688 ms 差得很多，但两次都编译了一个内核，差别不来自缓存；同样条件下其它几次运行在 100–260 ms 之间。

Vulkan 和 WebGPU（只读源码）

  • Vulkan。 burn 的 vulkan → cubecl 的 vulkan = ["wgpu-spirv"]，cubecl-wgpu 带上 spirv。设备是 Vulkan 适配器、扩展齐全时（register_vulkan_features），AutoCompiler 选 SpirvCompiler：IR 直接编成 SPIR-V 二进制，create_module 同样走透传，不经 naga。上面那个 SPIR-V 编译缓存只在这条路上生效，键是「适配器 vendor/device + 内核 + 构建 id」，换一次构建就失效。cubecl-wgpu 的 build.rs 在 macOS 上开了 spirv 却没设 VULKAN_SDK 时直接 panic（要经 MoltenVK）。本卷没有构建这条路。
  • WebGPU。 webgpu = ["wgpu"]，在 wasm 上 Auto 只有 BrowserWebGpu，编译器只有 WGSL；cubecl-server 的构建脚本把编译缓存限定在 not(target_family = "wasm")，注释是「the browser compiles its shaders itself, so it has no artifact to store」。WGSL 的写法和上面 rt-wgsl 的输出一样，只是由浏览器而不是 wgpu-hal 翻译。
❗IMPORTANT
动脑筋：换一种写法，要不要重新调优？

rt-tune 先在 Device::metal(DeviceKind::DefaultDevice) 上跑过 1024² 的 f32 matmul，autotune 结果已经写进 target/environment/default.db。下一个进程改用 Device::wgpu(DeviceKind::DefaultDevice)（同一块 GPU，同样是 wgpu<msl>），跑同一个 matmul。这次的第一次调用会启动几个 MatmulEntry 内核：1 个，还是几十个？答案在本页末尾。

动脑筋答案

几十个。依次是 Device::metal（缓存已有）、Device::wgpu 第一次、Device::wgpu 第二次，三个进程：

device: Device<Cube(Wgpu(WgpuDevice { kind: DefaultDevice, backend: Metal }))>
1024x1024 first:    120.127ms  launches=1 [MatmulEntry=1, ~compile(integrate)=1]
device: Device<Cube(Wgpu(WgpuDevice { kind: DefaultDevice, backend: Auto }))>
1024x1024 first:    778.378ms  launches=42 [MatmulEntry=41, UnaryNumeric=1, ~compile(integrate)=7]
device: Device<Cube(Wgpu(WgpuDevice { kind: DefaultDevice, backend: Auto }))>
1024x1024 first:    104.427ms  launches=1 [MatmulEntry=1, ~compile(integrate)=1]

数据库里多出一个命名空间：

autotune/0.11.3/device-100-0-wgpu<msl>/burn_cubecl-kernel-matmul-tune-base|1
autotune/0.11.3/device-100-16384-wgpu<msl>/burn_cubecl-kernel-matmul-tune-base|3
throughput/0.11.3/probe-v2/wgpu-msl-_spirv-0-0_Apple-M2-Max_mem4294967292_par0|4
  • autotune 的表名里有 CubeTuneId 的 Display 写出的 device-{type_id}-{index_id}-{runtime}。Auto 设备的 index 是 0，Metal 设备是 16384（图形 API 编在 index 的高位里，所以这两个 DeviceId 不相等）。同一块 GPU，两张表，Device::wgpu 那张是空的，于是 41 次候选测量、7 次编译。
  • 没有 LaunchOverhead 之类的测速内核：吞吐量缓存的命名空间按运行时名、适配器名和显存大小命名，不含设备 id，两种写法共用。
  • 第三个进程只剩 1 个：Device::wgpu 自己的表也有了。实际项目里固定用一种写法构造设备，能少一轮调优。

小结

  • burn 的 metal = cubecl 的 wgpu-msl：wgpu 运行时加 cubecl-cpp 的 MSL 编译器；原生 cubecl-metal（metal-native）burn 不用。后端名 wgpu<msl> 反映的是运行期是否真的用上了 MSL。
  • macOS 上 Device::wgpu 和 Device::metal 都落到 Metal API 和同一块 GPU，但设备 id 不同：两套服务线程、内存池、管线表和 autotune 表；Auto 在 MSL 不可用时可以退回 WGSL，Metal 不行。
  • MSL 是否可用在建服务端时决定：GPU 家族 Metal3 / Apple7 / Mac2 加一段只有 MSL 3.2 能编的 lambda 金丝雀。
  • 新内核的路径：KernelId 查表未中 → define_kernel 得 IR → MslCompiler 出 MSL 文本 → wgpu 透传 → Metal newLibraryWithSource → 计算管线，放回进程内的表。
  • 编译产物在 wgpu<msl> 上不持久化，compilation.cache = true 只对 SPIR-V（Vulkan）和原生运行时有用；每个新进程的第一次调用都要重新编译。
  • 只开 wgpu 时同一块 GPU 走 WGSL + naga，没有 CMMA，2048² matmul 慢三倍多。

💬 关联讨论

发起讨论
加载中...