CUDA（只读源码）

主线位置：两条主线换一个 engine。直接计算主线上，&mut engine 从 CpuEngine 换成 CudaEngine::open(..) 打开的值，add_read、dot_general_read 这些方法落到 CUDA 内核和 cuTENSOR；traced 主线上，runtime.prepare(&program, &reads, &[engine.target()]) 的目标换成 CUDA engine，run 时运行时调的是同一批方法。本机没有 NVIDIA GPU，也没有 CUDA 驱动、NVRTC 和任何 NVIDIA 库。 本页对内核编译、启动和厂商库调用的描述全部来自源码，没有运行过。实际跑过的只有不碰设备的部分：离线编译开了 cuda 的各 crate，调用设备探测入口看它在没有驱动时怎么失败，打印能力表和内核目录，以及让内核目录生成一个融合内核的源码。

flowchart TB
  eng["CudaEngine"] --> rt["CudaRuntime"]
  rt --> t4r["t4a-cuda-runtime"]
  eng --> disp["预备好的分派器"]
  disp --> km["KernelModules"]
  km --> cat["t4a-cuda-kernels"]
  km --> nvrtc["NVRTC 编译"]
  nvrtc --> cache["产物缓存"]
  cache --> disk["缓存目录"]
  eng --> ct["cuTENSOR 收缩"]
  la["tenferro-linalg"] --> sol["cuSOLVER cuBLAS"]
  ff["tenferro-fft"] --> cufft["cuFFT"]
  la --> eng
  ff --> eng

本页的实验 crate 是 exp-tenferro8/gpu-cuda（tenferro-gpu 0.8.0 开 cuda，外加 t4a-cuda-kernels 0.1.0 和 tenferro-tensor）、gpu-all（四个 crate 都开 cuda）和 gpu-feat（只给 tenferro-fft 开 cuda），release，M2 Max、macOS 27、rustc 1.99.0。

先看能不能编

tenferro-gpu 的 cuda feature 打开三个依赖：t4a-cuda-runtime 0.1.0、t4a-cuda-kernels 0.1.0 和 cudarc（0.19.10）。工作区给 cudarc 开的是 driver、runtime、nvrtc、dynamic-loading、cuda-12080，tenferro-gpu 再加 cublas、cutensor：编译时只要 CUDA 12.8 的绑定声明，不链接任何 NVIDIA 库，驱动和 NVRTC 都在第一次用到时 dlopen。t4a-cuda-kernels 的 CUDA 源码在 crate 里是 include_str! 进来的字符串，build = false，构建时不跑 nvcc。所以在没有 CUDA 工具链的 Mac 上，gpu-cuda 的 cargo build --offline --release 能编过，gpu-all 把 tenferro-gpu、-linalg、-einsum、-fft 都开 cuda 也能 cargo check 过。

tenferro-gpu 的 src/ 有 15536 行（不算测试文件），cuda/ 一个目录占 15127 行；lib.rs 里除了 #![deny(clippy::cast_ptr_alignment)] 这条 lint，所有模块都包在 #[cfg(feature = "cuda")] 里。主要文件：

内核目录是什么

0.7.1 的 CUDA 后端用 CubeCL：内核是 Rust 里的 #[cube] 函数，运行期由 CubeCL 生成 CUDA C++。0.8.0 换成了「kernel catalogue」：内核是一份 CUDA C++ 源码，单独发布成 t4a-cuda-kernels，它的 README 用一张表逐条写明每个入口是从 tenferro 的哪个 CubeCL 内核移植来的。它的包描述是「Runtime-independent CUDA kernels for tensor4all and external CUDA consumers」，lib.rs 开头写明了边界：

This crate owns kernel semantics, the C-compatible kernel ABI and checked
host-side planning. It does not depend on `t4a-cuda-runtime`, tenferro,
CubeCL or CubeK, never allocates device memory and never launches anything:
the caller supplies context, stream, buffers, module loading, workspace and
completion management.

所以一个「目录」有三样东西：cuda/ 下的 .cu 源码和共享头文件 abi.cuh（经 source::MODULES 和 source::ABI_HEADER 以字符串交出）；C ABI，入口名、参数顺序、按值传的布局结构 CLayout；Rust 一侧的 Entry 枚举和每个入口的 plan_* 函数，检查形状、步长、dtype，算出网格和块。它不编译、不分配、不启动，这些都是调用方的事。gpu-cuda catalog 打印目录本身：

ABI_REVISION = 11, Entry::ALL = 58
  arithmetic        4649 bytes   4 entries
  convert           8471 bytes   1 entries
  diagonal         15894 bytes   6 entries
  elementwise      24549 bytes   6 entries
  gather_scatter   32182 bytes   2 entries
  linalg_aux       48018 bytes  11 entries
  movement         34412 bytes   9 entries
  reduce           23319 bytes   2 entries
  rrqr             49112 bytes   8 entries
  validation        7400 bytes   5 entries
  window           21649 bytes   4 entries
ABI header 41904 bytes

58 个入口分在 11 个编译单元里；一个入口不只一个 __global__ 符号（.cu 里一共 94 个），因为秩不超过 8 的形状走把布局按值传进去的快路径，更高的秩走 _any 后缀的通用路径，从设备内存读布局。linalg_aux 和 rrqr 是 QR、LU、带列主元 QR 的辅助内核，给 tenferro-linalg 用。

tenferro 一侧不做校验。dispatch.rs 开头一句：「Every plan comes from the kernel crate's plan_* function — this module adds no validation of its own — and every argument order follows the documented ABI」。设备一侧的执行资源来自 t4a-cuda-runtime（包描述「Standalone CUDA execution runtime (streams, memory pools, transfers, NVRTC and kernel launch) for tensor4all」）：每个运行时一条真实的流、一个显式的内存池、按流排序的分配、锁页暂存、NVRTC 产物和模块缓存、带检查的内核启动。它也不认识 tenferro 的张量。两个 crate 互不依赖，tenferro-gpu 把它们接在一起。

编译一次，缓存两层

一个运算第一次用到某个入口时，engine 要准备这个入口所在的编译单元。KernelModules::prepare 先用 module_of(entry) 找到模块，跳过已经有的，然后对每个缺的模块：

        let capability = self.runtime.device_info().compute_capability();
        let arch = format!("sm_{}{}", capability.major, capability.minor);
        // ...
                let mut request = CompileRequest::new(module.name, compose(module));
                request.output = OutputKind::Cubin;
                request.arch = Some(arch.clone());
                let artifact = guard
                    .compile(&request)
                    .map_err(|error| crate::Error::runtime_state_source("cuda prepare", error))?;

compose 把 abi.cuh 拼到模块源码前面、去掉 #include 行，因为 NVRTC 吃的是内存里的字符串。目标架构取自设备的计算能力，产物是 cubin。按入口准备、按模块编译，一次运行只编它碰到的编译单元。

guard.compile 进的是 t4a-cuda-runtime 的产物缓存，有两层：

• 内存层：每个运行时一个，键是请求的哈希（源码 id、源码、头文件、编译选项、架构、输出种类），命中时还比对完整请求；默认上限 64 项。
• 磁盘层：由环境变量 TENFERRO_CUDA_ARTIFACT_CACHE 打开。default_config 读它，artifact_cache_dir 把空值和 none、off、memory 当成只用内存，其他值当目录，第一次写入时创建。注释给了它存在的理由：「a whole module is about 8 s per process on an A100, and a warm read is a few milliseconds」。

磁盘层的格式在 t4a-cuda-runtime 0.1.0 的 src/compile.rs（ArtifactCache::with_dir 的文档）：每个编译单元一个文件，名字是 v1-<source_id>-<128-bit digest>.<ptx|cubin>，摘要覆盖源码 id、源码、头文件、选项、架构、输出种类和 NVRTC 版本，所以改了一个模块只让那一个模块失效；写入先写临时文件再改名，并发读的进程看不到半个文件；目录读写失败只算一次未命中；旧文件不会自动删，「the directory is content-addressed, so an old entry is dead weight, not a hazard」。

两层缓存只省 NVRTC。拿到产物之后还要把模块装进上下文、绑定函数，这一步走运行时一个来回。PreparedDispatchers 让 engine 把准备好的模块集和绑定好的函数留着，后面的运算只在出现新入口时扩充：注释说这个来回「about a millisecond per module set; an engine that rebuilt one per operation paid it on every call」。提交 d8ddd68d（#2291）的说明记了实测：CopyExact 1.25 ms、Reduce 1.27 ms、Add 0.25 ms，「paying it per call is what made the U7 comparison look like a 15x regression」。

一次 add 在 engine 里

CudaEngine 的类型注释是「one runtime, one numerical submission path, one module cache」：一个 CudaRuntime、第一次收缩时才建的 cuTENSOR 句柄、扩展缓存、准备好的分派器、待回收的 prepared 运行，以及一个 broken 标记。它是 Send 不是 Sync，提交时要可变借用运行时；两个线程要共用一个 engine，得在外面加锁。和 CPU 一样，engine 由调用方持有、以 &mut 借出（见 engine 与 threadteam 池）。

CudaEngine::open 收的是 CudaDeviceId，写成 CudaEngine::open(CudaDeviceId::from_ordinal(0))；release notes 和用户指南写的 CudaEngine::open(ordinal) 是简写。它只做一件事：CudaRuntime::new(device_id)。上传下载不经过 engine，而是 upload_tensor(engine.runtime(), &cpu) / download_tensor(engine.runtime(), &gpu)，见 存储与所有权。

二元逐元素运算都走 elementwise_binary：

    fn elementwise_binary(
        &mut self,
        lhs: TensorRead<'_>,
        rhs: TensorRead<'_>,
        op: Elementwise,
        name: &'static str,
    ) -> crate::Result<Tensor> {
        let left = self.device_read(lhs, name)?;
        let right = self.device_read(rhs, name)?;
        // ...
        let kernel = kernel_dtype(dtype, name)?;
        let shape = left.shape().to_vec();
        let len = count_i64(&shape, name)?;
        let output = alloc_tensor(self.runtime(), dtype, &shape)?;
        // ...
        let mut dispatcher = self.dispatcher(&[entry])?;
        with_scalar!(dtype, |T| {
            dispatcher.arithmetic(
                entry,
                kernel,
                len,
                self.allocation_of::<T>(&output, name)?,
                self.allocation_of::<T>(&left, name)?,
                self.allocation_of::<T>(&right, name)?,
            )
        })?;
        self.runtime()
            .synchronize()
            .map_err(|error| crate::Error::runtime_state_source(name, error))?;
        Ok(output)
    }

device_read 读操作数：主机张量直接返回 Unsupported（「a host tensor would need an implicit transfer」），别的运行时的分配报错，借来的视图先物化成一份紧凑副本（engine 的 staged_reads 计数加一）。然后把 tenferro 的 dtype 翻成内核目录的 dtype，按流分配输出，取分派器，启动，最后 synchronize。两个操作数形状不同也返回 Unsupported。

最后那个 synchronize 是 CudaRuntime::synchronize，等流上排着的所有工作做完。structural.rs 里这样的调用有 45 处，copy_to_fresh 旁边的注释说明了用意：「The returned tensor must be complete when the caller sees it」。也就是说，直接调用的每个运算返回时，设备上的活已经干完了。用户指南 engines-and-transfers.md 的说法与此不同：「a chain of dependent operations is just a chain of calls, and one flush or synchronize covers them」，还说 flush 和 synchronize「Neither is implicit」；而 tenferro-gpu 里没有叫 flush 的方法。

prepared 程序的运行多一层。运行时调的还是这些方法（ExecutionEngine 以 Engine 为父 trait），运行结束把这次运行交给 retire_run：先放进队列，超过 MAX_PENDING_RUNS（16）才排空最老的一个；ExecutionEngine::synchronize 等流并排空整个队列，Drop 尽力同步一次。prepare / run 本身见 执行与缓存。

能做什么写在一张手写的表里：cuda_capabilities()，注释说它「is written down rather than derived from the kernel catalogue's accepts list」，由 tests/seam_capability.rs 对着 engine 的实际行为核对。它是 const fn，不碰设备，本机能打印：

cuda_capabilities(): 142 rows, Native 118, Unsupported 24
ops (31): Add Sub Mul Neg Conj Div Rem Abs Sign Maximum Minimum Compare Select Clamp Exp Log Sin Cos Tanh Sqrt Rsqrt Pow Expm1 Log1p Erf ReduceSum ReduceSumSquares ReduceProd ReduceMax ReduceMin DotGeneral

24 个 Unsupported 是复数的超越函数（Exp、Log、Sin、Cos、Tanh、Sqrt、Rsqrt、Pow、Expm1、Log1p 的 C32 / C64）、Bool 和复数的 Select、Bool 的 Compare。Erf 只有 F32 / F64。这张表只列核心算子；激活、softmax、归一化在 0.8.0 是由这些算子拼出来的组合运算（见 算子级 AD），在 CUDA 上能不能跑，取决于拼出来的每个算子在不在表里。

融合：零件在，没接上

内核目录还带一个融合内核的生成器 t4a_cuda_kernels::fusion：给一个 dtype、输入个数、按拓扑序排好的指令表和输出，emit_fusion 生成一段自包含的 CUDA C++ 源码和一个产物键。它同样不编译。gpu-cuda fusion 让它生成 exp(x * y) + x 的 f32 内核：

source 42489 bytes (ABI header 41904 bytes)
--- tail from first extern "C" ---
extern "C" __global__ void t4a_fused_elementwise(float* out0, const float* in0, const float* in1, long long len) {
  const long long index = static_cast<long long>(blockIdx.x) * blockDim.x + threadIdx.x;
  if (index >= len) {
    return;
  }
  const float v0 = in0[index];
  const float v1 = in1[index];
  const float v2 = t4a_cuda_mul(v0, v1);
  const float v3 = t4a_cuda_exp(v2);
  const float v4 = t4a_cuda_add(v3, v0);
  out0[index] = v4;
}

manifest = FusionManifest { entry: "t4a_fused_elementwise", abi_revision: 11, dtype: F32, inputs: 2, outputs: 1, block: 256, index_width: 64 }
key material 42605 bytes; sm_70 vs sm_80 equal: false
plan_fused(1) = Ok(Launch { grid: [1, 1, 1], block: [256, 1, 1], shared_bytes: 0, variant: Fast })
plan_fused(1048576) = Ok(Launch { grid: [4096, 1, 1], block: [256, 1, 1], shared_bytes: 0, variant: Fast })
plan_fused(0) = Ok(Launch { grid: [0, 0, 0], block: [256, 1, 1], shared_bytes: 0, variant: Fast })
complex exp: Err(FusionEligibility)

几点可以从输出里直接读到：源码的 98% 以上是嵌进去的 ABI 头文件，真正的内核十几行；元素个数 len 是运行期参数，形状和步长都不进源码，所以同一个融合计划对任意长度只编译一次；产物键包含目标架构，换一张卡就是另一个键；复数只收加减乘除这类代数运算，exp 被拒绝。

tenferro-gpu 有对应的零件 PreparedFusion::prepare：生成、编译、绑定，再按 len 启动。但 engine 并没有用它。impl TensorFusion for CudaEngine {} 是空的，execute_elementwise_fusion 用 trait 的默认实现，返回 Ok(None)，运行时于是在融合区域里逐条执行。整个工作区里除了它自己的测试，没有地方调 PreparedFusion。CPU 上 traced 程序的融合见 逐元素、SIMD 与融合；在 CUDA 上，v0.8.0 不融合。

厂商库：谁加载，缺了报什么

三处都用 libloading，按顺序试每个候选路径，全失败时把每个路径和它的 dlopen 错误拼进一个类型化错误：cuTENSOR 是 GpuError::Unavailable { detail }，cuSOLVER 是 Error::io_source("cuda_linalg", ..)，cuFFT 是 CudaFftError。没有一处回退到别的实现或 CPU。默认库名全是 Linux 的 .so。细节上有一处不同：cuSOLVER 和 cuFFT 的文档写着「A failed load is not cached」，环境变量在第一次成功之前每次都重读；cuTENSOR 的 load_library 把 Result 整个放进 OnceLock，第一次失败之后，同一进程里再改 TENFERRO_CUTENSOR_PATH 也没有用。

dot_general 一律是 cuTENSOR 收缩：submit_dot 把两边的模式、形状、步长写进描述符，交给 Cutensor::contract，计划按「dtype、模式、形状、步长」这个签名缓存，同一签名第二次调用不再做任何 cuTENSOR 规划。vendor.rs 里另有一个 cuBLAS 适配器 Cublas::gemm，能发 cublas?gemmStridedBatched，对外导出成 tenferro_gpu::cuda::Cublas，但工作区里只有 tests/seam_vendor.rs 调它，没有任何运算走这条路。capability.rs 里 dot_native 的注释「A borrowed operand reaches cuBLAS through a compact copy」说的也是 cuBLAS，和 submit_dot 的实际路径不符。

线性代数在 tenferro-linalg/src/gpu/，为 CudaEngine 实现：分解走 cuSOLVER，QR、LU、RRQR 的辅助步骤是内核目录里 linalg_aux、rrqr 两个模块的启动，cuSOLVER、cuBLAS 经裸会话在同一条流上调用。三个运算直接返回 Unsupported：full_piv_lu、full_piv_lu_solve（「complete-pivoting LU is not implemented on CUDA」），以及 eig，它的错误信息告诉你怎么办：「Download to CPU explicitly via backend.download_to_host(&gpu_tensor)? and then call CpuEngine::eig」，download_to_host 是 CudaEngine 的 TensorDeviceTransfer 方法。

solve 不再检查奇异（#2204，首次发布于 v0.8.0）：先 LU 分解、再回代，中间不读 cuSOLVER 的 info。提交 3abe5f5f 列了省掉的东西，「the epsilon-scaled diagonal singularity scan (extract_diagonal, magnitude, reshape, two reductions, a full stream synchronize, and two scalar downloads), the cuSOLVER getrf info download」，结论是「solve/lu_factor_solve therefore perform no per-call non-singularity check ... matching torch.linalg.solve」。所以在 CUDA 上对一个奇异矩阵 solve 不报错，返回的是 cuSOLVER 算出的什么值就是什么值；要检查奇异，用 lu_factor / lu，它们保留了 Singular 错误。CPU 一侧的行为见 线性代数。FFT 的 CUDA 路径见 FFT。

没有驱动时

库找不到报类型化错误，那连 CUDA 驱动都没有呢？cuda_devices 的签名是 Result<Vec<CudaDeviceInfo>, CudaDeviceError>，文档写着「Returns CudaDeviceError::Discovery for driver initialization, device enumeration, or device-name lookup failures」；lib.rs 的 crate 级示例和上游例子 cuda_quickstart 都是 cuda_devices()?。gpu-cuda 在 Mac 上分三次运行 gpu_available()、cuda_devices() 和 CudaEngine::open(CudaDeviceId::from_ordinal(0))。

下游内核和 ROCm

0.7.1 里，下游 crate 可以通过 cubecl-interop feature 拿到 CubeCL 会话，编译、启动自己的 #[cube] 内核。这条路随 CubeCL 一起删了。release notes 的「Known limitations」写得直接：「The CUDA raw session exposes the stream, allocations, copies and vendor handles, but no kernel compile-and-launch surface. Downstream #[cube] kernels, previously launched through the CubeCL session, have no replacement in this release」。

替代它的是 CudaEngine::with_raw 交出的 raw::Session：stream() 给出裸 CUstream，alloc_bytes、copy_bytes、upload_bytes、download_tensor 管字节，tensor / tensor_mut 借出张量的设备地址，resource 在 engine 的缓存里按类型放一个只初始化一次的对象（厂商句柄就放在这里）。tenferro-linalg、tenferro-fft 调 cuSOLVER、cuFFT 走的就是它。tenferro_gpu::cuda 还重新导出了 cudarc，让下游的厂商调用用同一版绑定。会话里没有编译和启动的入口；公开的 KernelModules 只认内核目录的 Entry，装不进自己的源码。下游要跑自定义内核，只能自己用 NVRTC 编、自己在这条流上启动。

ROCm 只是一个名字。tenferro-gpu 的 rocm = [] 什么都不打开，tenferro-einsum 和 tenferro-fft 的 feature 表写的是「Placeholder; HIP/ROCm is not implemented」。

一个 feature 陷阱还在

怎么读这一卷按源码推断过：tenferro-ad 的 EagerExtensionBackendKind 的 Cuda 变体跟着 tenferro-ad/cuda 走，开了 autodiff 的运算族 crate 对它 match，Cuda 那一支只在它们自己的 cuda 下编译。gpu-feat 只给 tenferro-fft 开 cuda（它会打开 tenferro-ad/cuda），tenferro-linalg、tenferro-einsum 只开 autodiff，cargo check --offline 以 101 退出：

error[E0004]: non-exhaustive patterns: `tenferro_ad::extension::EagerExtensionBackendKind::Cuda` not covered
   --> .../tenferro-einsum-0.8.0/src/eager_ad.rs:214:11
    |
214 |     match backend_kind {
    |           ^^^^^^^^^^^^ pattern `tenferro_ad::extension::EagerExtensionBackendKind::Cuda` not covered
...
error: could not compile `tenferro-einsum` (lib) due to 1 previous error
...
error[E0004]: non-exhaustive patterns: `tenferro_ad::extension::EagerExtensionBackendKind::Cuda` not covered
   --> .../tenferro-linalg-0.8.0/src/eager_ext.rs:982:11

和 0.7.1 的 WebGpu 版本同一个形状：一个 crate 打开 GPU feature，cargo 的 feature 统一把枚举变体加给了所有人，没打开的 crate 的 match 就不完整。用 CUDA 时，依赖图里每个开了 autodiff 的运算族 crate 都要一起开 cuda。gpu-all 那样四个一起开，就能编过。

动脑筋答案

== avail
gpu_available() = false
exit=0
== devices

thread 'main' (9602913) panicked at .../cudarc-0.19.10/src/lib.rs:200:5:
Unable to dynamically load the "cuda" shared library - searched for library names: ["libcuda.dylib", "libcuda64.dylib", "libcuda64_12.dylib", ...]. Ensure that `LD_LIBRARY_PATH` has the correct path to the installed library. If the shared library is present on the system under a different name than one of those listed above, please open a GitHub issue.
note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace
exit=101
== open

thread 'main' (9602919) panicked at .../cudarc-0.19.10/src/lib.rs:200:5:
Unable to dynamically load the "cuda" shared library - ...
exit=101

• gpu_available() 是 false，不 panic。gpu_available 调的是 t4a-cuda-runtime 的 device_count()，而 t4a-cuda-runtime 0.1.0 的 src/context.rs 里，init() 先用 is_culib_present() 探测驱动库，不在就返回 Error::DriverUnavailable（「libcuda could not be loaded」），gpu_available 把错误当成 false。设了 TENFERRO_REQUIRE_GPU=1 时它故意 panic（「TENFERRO_REQUIRE_GPU is set but no CUDA device is usable」），那是给 GPU 测试机用的。
• cuda_devices() 既不是 Ok(vec![]) 也不是 Err，而是在 cudarc 里 panic，进程以 101 退出。它不经过 t4a-cuda-runtime：discover_with 的第一步是 CudaDriverApi::initialize，直接调 cudarc::driver::result::init()；cudarc 开了 dynamic-loading，第一次调驱动函数时才加载 libcuda，加载失败就在它自己的 lib.rs 里 panic，tenferro 没机会把它变成 Discovery。
• CudaEngine::open 也 panic，RUST_BACKTRACE=1 显示它的路径是 CudaEngine::open → CudaRuntime::new → CudaRuntime::with_config → cuda_devices → cudarc::driver::result::init → cudarc::panic_no_lib_found。with_config 先调 cuda_devices() 查设备信息，再调 t4a-cuda-runtime 的 Runtime::open；后者本来会返回类型化的 DriverUnavailable，但轮不到它。
• 上游知道 cudarc 会 panic：CudaDeviceInfo 几个方法的文档示例都写着「cudarc panics when the CUDA driver library is absent, so the call is guarded (the same pattern gpu_available uses)」，用 catch_unwind(cuda_devices) 包起来。但 cuda_devices 自己的文档、crate 级示例和 cuda_quickstart 都直接 ?。在可能没有 CUDA 的机器上，先调 gpu_available()，为 true 再 open。

小结

• CudaEngine 是一个 t4a-cuda-runtime 运行时加一套缓存；内核来自 t4a-cuda-kernels 的内核目录：58 个入口、11 个 CUDA C++ 编译单元、C ABI 和 Rust 侧的 plan_* 校验，目录本身不编译、不启动。
• 内核按模块用 NVRTC 编成 cubin，目标是设备的计算能力；内存层缓存在运行时里，TENFERRO_CUDA_ARTIFACT_CACHE 打开按内容寻址的磁盘层；engine 留着准备好的模块和函数，避免每次运算一个来回。
• 直接调用的每个运算返回前都 synchronize；prepared 运行最多攒 16 个再回收。用户指南说的「一次 synchronize 覆盖一串运算」和 flush 方法，在源码里对不上。
• dot_general 只走 cuTENSOR；Cublas::gemm 有但没人调，#2193 的批量 GEMM 快路径在发布前随 CubeCL 删掉了，release notes 还在说它。融合内核的生成器和 PreparedFusion 都在，engine 的 TensorFusion 是空实现，v0.8.0 的 CUDA 不融合。
• cuTENSOR、cuSOLVER、cuFFT 运行期加载，可用环境变量覆盖，缺库报类型化错误、不回退；eig、full_piv_lu、full_piv_lu_solve 返回 Unsupported；solve 不检查奇异。
• 没有驱动时只有 gpu_available() 安全，cuda_devices() 和 CudaEngine::open 在 cudarc 里 panic。下游 #[cube] 内核没有替代入口；ROCm 是空 feature；只给一个运算族 crate 开 cuda 会编不过。