feat(accel): version floors, a probe channel, constrained device globs, and a CUDA example that takes two routes

.agents/docs/2026-09-05-multi-device-ecosystem-design.md

@@ -1180,3 +1180,97 @@ lane 结论一致;只跑模拟器,等于把「lavapipe 说对了」当成「它
   第三档的「CUDA 只能上真卡」删除,代之以 §5 的矩阵;
 - **§12** 的阶段表 —— 「未做」的理由全部作废,代之以 §5 的 lane 与 §9 的判据;
 - **§15.8** —— 两条「缺授权」的结论作废,代之以 §4.2 与 §3.1。
+
+---
+
+## 12. 实施后自我 review(2026-09-05)
+
+方案写完之后由实施推翻或补上的地方。**只记与本文所写不同的**;相符的部分见
+实施计划表的判据列。
+
+### 12.1 被实施推翻的
+
+| 本文写的 | 实测 | 处置 |
+|---|---|---|
+| C-1「设备目标原语」是新的 `[[target]] kind = "device"` | `mcpp::action` 已有四种角色,artifact 的产物不进链接、object 的进链接,且 ninja 按路径连边 —— 「不参与常规链接、被某条边消费的产物」**就是** artifact 角色 | 不新增 target kind。再加一种是同一个决定写第二遍 |
+| C-5 的驱动取数由核心完成 | 仓库自带的 `test_runtime_contract` 禁止「厂商词 + 探针启动」在 `src/` 共现,抓住了写下的 `nvidia-smi` 调用 | 关系留在核心并单测,**取数改由声明抵达**;`doctor` 的整节与 `mcpp.toolchain.devicehost` 一并删除 |
+| C9(RDC 真机)在本机可测 | nvcc 路线**两端同时被挡**:12.9 满足驱动而不满足 C 库(C23 `cospi`/`rsqrt` 的 `noexcept` 冲突),13.3 满足 C 库而不满足驱动(要 ≥13.0,本机 12.4) | 判据退回 e2e 607 所测的**通用链式 action**;真机 RDC 留给有 13.x 驱动的机器,并在示例 README 里写明为什么 |
+| ① 已完成 | 13.x 的 `cuda-nvcc` **装完不能用**:nvcc 用 `$(TOP)/nvvm/bin/cicc` 找后端,而 13.x 把 `nvvm/` 与 `crt/` 拆成了独立包=独立载荷根 | xim #760。修法的形状:**无条件写链接**,不要求嵌套安装成功 —— 它失败无声且两种拼法都不可靠 |
+| T1.5「lavapipe 载荷」是新增包 | 现有 `xim:mesa` 载荷里**只有 RADV**,没有 `libvulkan_lvp.so`,也没有 rusticl | 仍待做,且不是新增包而是**重打 mesa 载荷**(`-Dvulkan-drivers=…,swrast` + `-Dgallium-rusticl=true`)。T1.4 同理 |
+
+### 12.2 实施自己造出来又修掉的
+
+- **「没有加速器」被写成了显示用的 `(none)`。** `accel_str` 为空集打印 `(none)`
+  是给 ABI 标签读的;`resolvedAccel` 把这个拼法当值传了出去,于是
+  `MCPP_ACCEL=(none)` 到达**每一个从未提过加速器的工程**,而指纹里
+  `if (!accel.empty())` 恒真。e2e 605 第四段标题写着「变量与 layer 都清空」
+  却只测了 layer —— 这就是它逃过套件的原因。判据只能靠构建程序**写文件**取得:
+  它的 stdout 只在非零退出时才打印。
+- **设备源的映射按裸包名索引。** 同一张图里两个包可以同名不同命名空间;
+  改按包根索引。合入前重读 diff 时发现,没有测试覆盖它。
+
+### 12.3 生态级 review:这一轮之后,一个消费者看到的是什么
+
+一个要用 GPU 的工程现在写三样东西,各自答给不同的所有者:
+
+```toml
+[xlings.workspace]                  # 载荷:工程自己选版本
+"xim:cuda-nvcc"   = "12.9.86"
+
+[dependencies.compat]               # 机器:驱动由机器决定,包只负责够到它
+cuda-driver = "2026.09.05"
+cublas      = "12.9.1.4"            # 算子库:载荷 + 构建面,两个仓库各管一半
+
+[build]                             # 轴:写一次,规则包据此推导自己的开关
+accel = "cuda12.9+{sm_89} ptx>=89"
+```
+
+**核心不认识其中任何一个厂商名字**,这条由 `test_core_vendor_probes` 在剥掉
+注释的源码上执行,并自带分母。
+
+三条已被生态执行的规则在本轮各验证一次:
+GPU 索引包不自己探测宿主(委托 xim sentinel);链宿主 `libcudart` 会被闭包
+校验拒掉(所以 `compat.cudart` farm 的是载荷而不是宿主);打包后的
+`runtime.artifacts` 是封闭白名单(新增产物字段两个读取器都要查)。
+
+### 12.4 仍然没有做的,以及理由
+
+| 项 | 理由 |
+|---|---|
+| T1.2 `llvm-offload` | `dpcpp@7.1.0` 载荷自带全套 offload 工具,需要 RDC 的工程可用它;独立包仍待做 |
+| T1.4 / T1.5(pocl / lavapipe)与 T4.2 | 需要**重打 mesa 载荷**(见 12.1 末行),或新建 pocl 源码构建配方。两者都是多小时的载荷工程 |
+| T1.8/T1.9/T1.10(chipstar / adaptivecpp / hip) | 依赖 T1.2/T1.4 |
+| T4.3 规则包进索引 | 依赖 ③ —— 描述符指向 mcpp 的**源码 tarball**(`grpcgen` 同形),tag 不存在则算不出 sha256。规则包已改名到 `mcpplibs` 命名空间,就是为了让它可被引用而不是被复制 |
+| ⑤ 九个框架 | 依赖 ④ 的规则包条目。`ggml-org.llamacpp` 与 `opencv.opencv` 已在索引里,多后端是改**它们各自的 `-m` 仓库**而不是索引条目。T5.1 已做到「链路全通、卡在载荷矩阵」—— 见 12.5 |
+| T2.6 的端到端判据 | `accel` 已是 `pack::AbiTag` 第四维并进指纹;「`.a` 随包传播」还缺一条跨包的判据 |
+
+### 12.5 T5.1 作为 gate 的实际读数
+
+**它兑现了 gate 的作用**:第一个真实框架就暴露了 C-6 的引擎缺口 —— object 角色的
+action **只**挂到可执行/共享库/测试上,而 llama.cpp 的 CUDA 后端是 305 个 `.cu`
+挂在 `kind = "lib"` 上,于是每个 action 都被丢弃、只留一条警告,**构建成功**并产出
+一个不含设备码的归档。修好并有判据(e2e 608 断言 `ar t` 的成员表 —— 空档案也会
+成功退出)。
+
+**链路本身全通,实测到 48 个设备目标**:`[build] accel` → 带 `accel` 的 glob →
+`MCPP_DEVICE_SOURCES` → 规则包 → `mcpp::action` → 归档 → 链接。
+
+**挡住的是一个四维载荷矩阵,四条边没有一条是 mcpp 的:**
+
+| 组合 | 读数 |
+|---|---|
+| CCCL 2.x(12.9 线)+ clang | `cub::LoadDirectWarpStriped` 少一个四参重载 |
+| CCCL 3.3(13.3 线)+ clang | 同一个调用,候选是三参与五参 |
+| CCCL 3.2(13.2 线)+ clang | 换成 **libcu++ 编不动**:`string_view` 的推导指引只允许 `__host__ __device__`;`block_load.cuh` 要 placement new |
+| 任一 CCCL + nvcc | 12.9 撞 glibc 2.44 的 C23 `cospi`;13.3 撞驱动 12.4 |
+
+⭐ **这不是「没做完」,是「本机构造上无解」**,与 §12.1 里 C9 那条同一性质。
+需要的是一台驱动 ≥ 13.0 的机器(nvcc 13.3 路线),或一个 ggml 与 CCCL 版本匹配的
+上游 checkpoint。
+
+⭐ 顺带三条通用读数,都写进了规则包与文档:
+**layer 不能选择依赖**(依赖挂 feature,源文件挂 accel 轴);
+**设备编译必须指名 CCCL 载荷**否则命中 `/usr/include/cub`(与 §12.1 的
+`cuda_runtime.h` 同一形状,第三次);
+**clang 路线要带 `-D_ALLOW_UNSUPPORTED_LIBCPP`**,因为 NVIDIA 那条 `libc++ is not
+supported` 的守卫看的是 `__CUDACC__`,而 clang 编 CUDA 时自己就定义它。