diff --git a/.agents/docs/2026-09-05-multi-device-ecosystem-design.md b/.agents/docs/2026-09-05-multi-device-ecosystem-design.md new file mode 100644 index 00000000..8ca935d2 --- /dev/null +++ b/.agents/docs/2026-09-05-multi-device-ecosystem-design.md @@ -0,0 +1,1182 @@ +# 多设备加速生态:载荷、适配面、验证 lane 与框架验证 + +2026-09-05。本文取代 `2026-09-05-accelerator-support-design.md` §8.5 / §12 / §15.8 +三处的结论,并把范围从 CUDA 扩到全部主流设备后端。 + +已落地的引擎侧能力见前文;本文只写**尚未做的部分**,以及做它们的确切路径。 + +--- + +## 0. 前一版被实测推翻的六处 + +| # | 前一版说 | 实测 | 出处 | +|---|---|---|---| +| 1 | 「CUDA 工具链载荷本机无法真实验证」 | **载荷是 xim 的事,不是本机的事。** NVIDIA 的 redist manifest 给出每个组件的 URL/sha256/size/**上游版本号** | `redistrib_13.3.0.json` | +| 2 | 「`compat.cublas` 在索引的两种形态里都不成立」 | 说法错。第三种形态存在:**载荷进 xim,适配进 mcpp-index**,用 `xpm..deps` 接线 | 该机制我已在 `compat.cuda-runtime` 用过一次 | +| 3 | 「`rules-cuda` 进索引要新建仓库」 | 错。`grpcgen` 就是 `mcpp = "*/rules/mcpp.toml"` 指向**已有仓库的子路径**;mcpp 自己的 release tarball 里已含 `examples/09-cuda-kernel/rules-cuda` | `pkgs/g/grpcgen.lua` | +| 4 | 「决定 4(capability 同名符号边界)不可实现」 | 不可实现的是「一刀拒」那个实现。加一个 manifest 声明即可,是设计决定 | 见 §6.4 | +| 5 | 「CUDA kernel 的执行只能上真卡」(§8.5 第三档) | 错。**chipStar v1.3.0**(CUDA/HIP → SPIR-V → OpenCL)配 **PoCL v7.2** 的 CPU 设备,这条链是活的 | 两者均在 2026-09 有发布/提交 | +| 6 | 「Metal 本机无 macOS」 | 理由不成立。仓库有四条 macOS CI | `ci-macos`、`ci-macos-e2e`、`bootstrap-macos`、`macOS ARM64 — xlings LLVM end-to-end` | + +⭐ 六条里有五条的形状相同:**我把「我手边没有」写成了「做不到」**,而 xlings 生态 +的存在意义正是消除「手边有没有」这个变量。 + +--- + +## 0.1 ⚠️⚠️ 已发布的示例违反了本文的规则,两处,必须先修 + +`examples/09-cuda-kernel`(2026.9.5.1 已发布)是**教用户怎么做**的东西, +而它教的两件事都是错的。 + +### 0.1.1 用了已被取代的 `[xlings] deps` + +```toml +[xlings] +deps = [{ linux = "libcuda-host-link" }] # ⚠️ 2026.9.3.1 起已被取代 +``` + +`[xlings.workspace]` 才是**那一张表**(#548)。`deps` 仍被接受,所以不报错 —— +**这正是它危险的地方**:示例不会红,而抄它的人写下的是过时形式。 + +⚠️ 且这个键在文档里**到处都是**:英文侧 `docs/05` ×5、`07` ×2、`13` ×2、 +`15` ×1、`17` ×3、`18` ×1;中文侧 `05` ×3、`07` ×2、`13` ×2、`17` ×2、`18` ×1。 +⇒ **从全部用户文档中删除,不留「也可以这么写」** —— 留着就是留一个误用入口。 +(`[xlings].subos` 与 `[xlings.workspace]` 不受影响,它们是当前形式。) + +### 0.1.2 手写了 host 路径 + +```toml +[build] +ldflags = ["-L/usr/local/cuda/lib64", "-L/usr/lib/x86_64-linux-gnu", + "-lcudart_static", "-lrt", "-lpthread", "-ldl"] +``` + +规则包里同样: + +```cpp +c.push_back("/usr/local/cuda/bin/nvcc"); +c.push_back("/usr/bin/nvcc"); +c.push_back("/usr/include/crt/host_config.h"); +``` + +⚠️⚠️ **示例通篇依赖宿主的 CUDA 工具包**,而本文 §1 的规则说 host 只留 +`libcuda.so.1`。**我写了一个示范违反自己规则的示例。** + +⇒ 正确形态,在 §3.1 的载荷落地后: + +```toml +[xlings.workspace] +"xim:cuda-nvcc" = "12.9.86" +"xim:cuda-cudart" = "12.9.79" + +[dependencies.compat] +cuda-driver = "2026.09.05" # 唯一的 host 面,由包持有 + +[build] +accel = "cuda12.9+{sm_89} ptx>=89" +# ldflags 里不出现任何绝对路径 +``` + +规则包侧:nvcc 与 `crt/host_config.h` 都从 `mcpp::xpkg_dir("xim","cuda-nvcc")` +得到,**一个 `/usr` 都不出现**。 + +⭐ 这两条排在 §8 的最前面 —— 它们不是新功能,是**已发布的错误示范**, +每多一天就多一批照着抄的人。 + +## 1. 一条规则,以及它的判据 + +> **凡是能进 xim / mcpp-index 的,一律进;host 只保留物理上无法再分发的那一层, +> 并且那一层的边界必须被测量,不被假设。** + +判据不是「我们尽量了」,而是可以逐项核对的两问: + +1. **这个东西为什么不在载荷里?** 答案只允许是「许可禁止再分发」或 + 「与内核 ABI 锁步」。其余任何答案(体积、麻烦、本机已有)都不成立。 +2. **它的版本边界是被问出来的还是被假设的?** 每一条 host 依赖必须有一个 + **运行期可查的问法**,mcpp 用那个问法回答,而不是内置一张表。 + +⭐ 这是同一条原则的第三次应用: +- `crt/host_config.h` → 宿主编译器上界(已落地) +- `nvcc --dryrun` → 设备编译器能否够到自己的后端(已落地) +- `cuDriverGetVersion` → **驱动能跑多新的工具包**(本文新增,见 §2.1) + +--- + +## 2. host 必须面的穷举 + +逐后端列出「哪一层必须来自 host」,以及为什么不能进载荷。 + +| 后端 | host 必须面 | 为什么不可再分发 | mcpp 怎么问它 | +|---|---|---|---| +| **CUDA** | `libcuda.so.1`(驱动用户态库) | NVIDIA 驱动 EULA 禁止第三方再分发;且与内核模块 ABI 锁步 | `cuDriverGetVersion` | +| **ROCm/HIP** | `libamdhip64` 之下的 KFD/内核驱动接口 | 内核侧 | `hipRuntimeGetVersion` | +| **Vulkan** | ICD(`/usr/share/vulkan/icd.d/*.json` + 厂商 `.so`) | 厂商驱动;**但软件 ICD(lavapipe)可以进载荷** | `vkEnumeratePhysicalDevices` | +| **OpenCL** | ICD(`/etc/OpenCL/vendors/*.icd`) | 同上;**PoCL 可以进载荷** | `clGetPlatformIDs` | +| **Level Zero** | `libze_loader` + 厂商 driver | 厂商驱动 | `zeInit` | +| **Metal** | 系统框架 | 属于 macOS | 平台判定 | + +其余一切 —— nvcc、cudart、nvrtc、cuBLAS、hipcc、DPC++、AdaptiveCpp、 +PoCL、lavapipe、SPIR-V 工具、chipStar —— **全部进 xim**。 + +### 2.1 ⭐⭐ 驱动版本决定可用的工具包版本 —— 已实测,失败很晚 + +**实测(2026-09-05,本机 RTX 4080 / 驱动 550.144.03):** + +用 CUDA **13.3.33** 载荷编 `sm_89`,与用宿主 **12.0** 编同一份源码,同一台机器: + +| 用什么编 | 编译 | 链接 | 运行 | +|---|---|---|---| +| CUDA 13.3 载荷 | ✅ 干净 | ✅ 干净 | ❌ `runtime=13030 driver=12040`
`cudaMalloc: CUDA driver version is insufficient for CUDA runtime version` | +| 宿主 CUDA 12.0 | ✅ | ✅ | ✅ `result: 12 24 36 48` | + +⚠️⚠️ **编译与链接全干净,失败在第一次 `cudaMalloc`。** 这与 §1 那个「为 sm_90 编的 +库被 sm_86 消费」是同一类失败:晚、且消息不指向任何一次人为选择。 + +⇒ **处置:两条版本线并存,由 xim 的多版本管理承担。** 13.3 线与 12.9 线各自完整, +mcpp 在构建前问驱动要一个数,选能跑的那条。**这不是权变,是 xim 本来就是干这个的。** + +### 2.2 ⭐⭐ 载荷是一个闭包,而这个闭包要被**发现**,不能被猜 + +同一次实测暴露的第二件事:CUDA 13.x 把组件拆得比 12.x 细得多。 +只装 `cuda_nvcc`(30.1 MB)会得到: + +``` +sh: 1: .../bin/../nvvm/bin/cicc: not found +``` + +`cicc` 在 **`libnvvm`**(46.9 MB),`crt/host_config.h` 在 **`cuda_crt`**, +另有 `libnvptxcompiler`、`libnvfatbin`、`cuda_culibos`、`cccl`。 +一个**能编东西的** nvcc 是七个组件的闭包,不是一个包。 + +⭐⭐ **补齐它的办法不是抄一张目录表,而是用已经落地的那条 doctor 检查驱动它** —— +`nvcc --dryrun` 报出第一个解析不到的裸名,装上带它的组件,再问一次。 +这与 §15.7 是同一个机制:**那条检查原本用来抓坏掉的宿主,它同样抓不完整的载荷。** + +### 2.3 ⭐ 载荷带着自己的 `nvcc.profile`,因此没有沙箱那个病 + +Debian 把 `nvcc.profile` 做成指向 `/etc` 的符号链接,沙箱替换 `/etc` 就断 +(前文 §15.7)。**redist 载荷把 `nvcc.profile` 放在自己的 `bin/` 里** —— +解到任意前缀都自定位。实测 `nvcc --dryrun` 声明的 PATH 指向载荷自己的目录。 + +⭐ 顺带量到:**载荷自带的 `host_config.h` 上界是 gcc ≤ 15 / clang ≤ 21**, +而宿主那个 CUDA 12.0 是 gcc ≤ 12。⇒ **载荷不只是「不依赖 host」,它的宿主编译器 +上界还宽得多** —— 这是「进生态」的直接收益,不是理念收益。 + +## 3. xim 载荷表(全部取当前最新,附实数) + +版本号一律用上游的。日期版本只留给「内容是本生态自己写的、没有上游产物」的适配器 +(与 `compat.glx-runtime` / `compat.vulkan-runtime` 同规)。 + +### 3.0 ⭐⭐ 载荷边界由 manifest 的 `license` 字段决定,不由我判断 + +CUDA 13.3.0 的 linux-x86_64 组件共 **43 个、4.19 GB**,manifest 把它们分成三类许可, +而这三类的处置各不相同: + +| 许可 | 组件 | 合计 | 处置 | +|---|---|---|---| +| `CUDA Toolkit` | nvcc、cudart、nvrtc、nvjitlink、**cuda_gdb**、**sanitizer**、cuobjdump、nvdisasm、cuxxfilt、nvml、opencl、profiler_api、nvtx… | ~1.9 GB | **全部进 xim** | +| `NVIDIA SLA` | **nsight_systems 1074.8 MB**、**nsight_compute 336.5 MB** | 1.41 GB | **进 xim**,GLOBAL url 直指 NVIDIA 自己的 CDN | +| `NVIDIA Driver` | nvidia_driver 528.9 MB、cuda_compat | 0.53 GB | **唯一留在 host 的那层** | + +⭐ 第二行不是让步,但我原先的说法过绝对,已按 xim 的规范改正: + +- xim **推荐**官方二进制走 `xpm.source = "xlings-res"`,那要求 + `github.com/xlings-res/` 与 `gitcode.com/xlings-res/` 两腿都有同名 + release、资产字节一致 —— 那是 re-host; +- 但规范**明确允许**第三方 release 用 **URL template + per-arch sha256** + (`source = "https://.../${version}/...${arch}..."`),索引里 247 个配方直连 + github 就是这一形态。 + +⇒ **CUDA 组件走 URL template 直连 NVIDIA 的 CDN**:manifest 已给出稳定 URL 与 +sha256,不产生再分发行为,SLA 与 Toolkit 两类许可都成立。 +⚠️ 代价是 CN 用户直连 NVIDIA CDN;若实测过慢,再单独决定是否把第①档 +(150 MB)镜像进 xlings-res —— **那是一个可以后加的优化,不是前置条件**。 + +⚠️ **「体积」不是任何一条的理由。** 去掉 driver 后 **3.66 GB 全部进 xim**, +拆成约 40 个可独立安装的组件包 —— 一次普通 CUDA 构建只装第①档的 **150 MB**, +要 profile 的人才装 nsight。**这正是组件级拆包存在的意义。** + +### 3.1 CUDA 载荷(全部进 xim,版本一律用 manifest 里的上游版本) + +**⓪ 编译器本身 —— 已在 xim,不是新包** + +| xim 包 | 版本 | 说明 | +|---|---|---| +| `xim:llvm` | **22.1.8**(已装) | ⭐ **加速器路线的默认编译器**(§6.3.5)。它同时是宿主编译器、CUDA 设备编译器(`-x cuda`)、SPIR-V 前端 | +| `xim:gcc` | 13.3.0 / 15.1.0 / 16.1.0(已装) | 岛形态里设备 TU 的 `-ccbin` 备选;普通 C++ 工程的默认不变 | + +⭐ 主路线的编译器**一个新包都不用加** —— 它已经是 mcpp 生态的一等公民。 +这正是「以 LLVM 为主」在工程上便宜的原因。 + +**① CUDA 编译与运行必需 —— 150 MB** + +| xim 包 | 13.3 线 | 12.9 线 | 大小(13.3) | +|---|---|---|---| +| `xim:cuda-nvcc` | 13.3.33 | 12.9.86 | 30.1 MB | +| `xim:cuda-cudart` | 13.3.29 | 12.9.79 | 1.5 MB | +| `xim:cuda-nvrtc` | 13.3.33 | 12.9.86 | 66.4 MB | +| `xim:libnvjitlink` | 13.3.33 | 12.9.86 | 53.5 MB | +| `xim:cuda-cccl` | — | 12.9.27 | 1.0 MB | +| `xim:cuda-nvtx` / `xim:cuda-profiler-api` | 13.3.29 / 13.3.27 | | 0.1 MB | + +两条线并存,由 §2.1 的驱动判据选。⚠️ 本机驱动只到 CUDA 12.4 ⇒ 走 12.9 线。 + +**② 算子库 —— 按 shared / static 拆开** + +宿主上实测的 shared 实体大小(这是消费者真正要装的): + +| 库 | shared | 对照:static 档案 | +|---|---|---| +| `libcublas.so` | 102.5 MB | `libcublas_static.a` 154.9 MB | +| `libcublasLt.so` | 484.2 MB | `libcublasLt_static.a` **823.6 MB** | +| `libcusolver.so` | 290.4 MB | | +| `libcusparse.so` | 243.4 MB | | +| `libcufft.so` | 147.0 MB | | +| `libcurand.so` | 92.2 MB | | + +⇒ `xim:libcublas` 拆成 `libcublas` / `libcublas-static` 两个包:默认装 shared, +需要静态链接的人显式装。**不是砍功能,是让消费者付自己那份。** + +**③ 调试与分析 —— 生态完备性要求,全部进** + +| xim 包 | 版本 | 大小 | 许可 | +|---|---|---|---| +| `xim:cuda-gdb` | 13.3.27 | 90.1 MB | CUDA Toolkit | +| `xim:cuda-sanitizer` | 13.3.27 | 10.2 MB | CUDA Toolkit(compute-sanitizer) | +| `xim:cuda-nvdisasm` / `cuobjdump` / `cuxxfilt` | 13.3.29 | 4.8 MB | CUDA Toolkit | +| `xim:nsight-systems` | 2026.1.3.243 | 1074.8 MB | NVIDIA SLA | +| `xim:nsight-compute` | 2026.2.0.7 | 336.5 MB | NVIDIA SLA | + +⭐ 少了这一档,mcpp 上的 CUDA 开发只能编不能查 —— **那才是偷工减料**。 +`mcpp` 侧接线:`mcpp run --profile nsys` / `--check sanitizer` 走既有的 +runner 与 action 机制,不新增专用命令。 + +### 3.2 SYCL / oneAPI + +| xim 包 | 版本 | 大小 | 来源 | +|---|---|---|---| +| `xim:dpcpp` | **7.1.0**(2026-09-02) | 207 MB(linux) | `intel/llvm` release `sycl_linux.tar.gz` | +| `xim:adaptivecpp` | **25.10.0** | 源码构建 | 无 release 资产 ⇒ 依赖 `xim:llvm` 构建 | + +⭐ **已核实:官方 Linux 资产已含 CUDA 与 HIP 后端。** v7.1.0 的构建配置是 +`--cuda --hip`,release notes 的测试矩阵里有 "NVIDIA CUDA BACKEND on NVIDIA +GeForce RTX 3090"。⇒ Linux 不需要自建。 + +⚠️ **Windows 不含**:同一份 release notes 写着 "HIP & CUDA plugins on Windows are +not being built"。Windows 上自建,而**版本号仍用上游 tag `7.1.0`** —— 见 §3.5。 + +### 3.3 CPU 设备与模拟器(这些也进载荷,不依赖 host) + +| xim 包 | 版本 | 提供什么 | +|---|---|---| +| `xim:pocl` | **7.2**(2026-09-04) | OpenCL **CPU 设备** | +| `xim:mesa-lavapipe` | 25.2.8+ | Vulkan **软件 ICD**(`PHYSICAL_DEVICE_TYPE_CPU`) | +| `xim:chipstar` | **1.3.0** | CUDA/HIP → SPIR-V → OpenCL/Level Zero | +| `xim:spirv-tools` / `xim:glslang` | 最新 | SPIR-V 汇编与 GLSL→SPIR-V | + +⭐⭐ **模拟器进载荷,是「无卡 CI 也能跑 kernel」这件事从偶然变成契约的唯一办法。** +本机现在能跑 lavapipe 是因为发行版恰好装了它;进 xim 之后,任何机器都能跑。 + +### 3.4 ROCm / HIP + +| xim 包 | 版本 | 说明 | +|---|---|---| +| `xim:hip-runtime` / `xim:hipcc` | ROCm **7.14.0**(2026-07-16) | ⭐ HIP 可以 `HIP_PLATFORM=nvidia` 走 CUDA 后端 ⇒ **本机 4080 就能验 HIP 前端** | + +### 3.5 ⚠️⚠️ 版本对齐是硬规则,自建也不例外 + +> **凡是有上游版本的,一律用上游版本;自己编的不改号。** + +- `xim:cuda-nvcc@13.3.33` —— manifest 里的 `version` 字段,逐字。 +- `xim:dpcpp@7.1.0` —— 上游 tag。**Windows 自建出来的也叫 7.1.0**,因为它就是 + 那份源码;构建差异记在描述符的构建元数据里,不体现在版本号上。 +- `xim:adaptivecpp@25.10.0` —— 上游 tag,尽管是源码构建。 +- `xim:pocl@7.2`、`xim:chipstar@1.3.0`、`xim:hip-runtime@7.14.0` —— 同理。 + +日期版本(`2026.09.05` 这种)**只留给内容是本生态自己写的、没有上游产物的适配器**, +与 `compat.glx-runtime` / `compat.vulkan-runtime` 同规。 + +⚠️ 这条已经被违反过一次:`compat.cuda-runtime` 的名字与内容不符(§4.1), +虽然版本形态是对的。 + +--- + +## 4. mcpp-index 表 + +xim 装载荷,mcpp-index 描述「怎么对它构建」。实测:xim-pkgindex 179+ 个包里 +**零个带 `mcpp` 块** ⇒ 分工是硬的,适配器必须住在 mcpp-index。 + +### 4.1 适配器(host 面,farm + `runtime.library_dirs`) + +| 包 | 状态 | 动作 | +|---|---|---| +| `compat.cuda-driver` | ⚠️ 现名 `compat.cuda-runtime`,**名字与内容不符** | **改名**。NVIDIA 词汇里 "CUDA Runtime" 专指 `libcudart`;本包 farm 的是驱动,其 `capabilities`/`provides` 已写作 `cuda.driver` | +| — | ⚠️ `repo` 字段指向 `openxlings/xim-pkgindex` | **改正**。同族两包指向被适配物的上游 | +| `compat.vulkan-icd` | 新增 | farm host ICD;缺失时回落 `xim:mesa-lavapipe` | +| `compat.opencl-icd` | 新增 | 同上,回落 `xim:pocl` | + +### 4.2 构建规则包 + +| 包 | 来源 | 说明 | +|---|---|---| +| `mcpplibs.rules-cuda` | mcpp release tarball 的 `examples/09-cuda-kernel/rules-cuda` | ⭐ **一个仓库都不用新建**,与 `grpcgen` 同形。搬进索引后示例改成消费索引里那一份,避免两份 | +| `mcpplibs.rules-hip` | 同一 tarball | hipcc 的 `--offload-arch` | +| `mcpplibs.rules-spirv` | 同一 tarball | `glslc` / `slangc` → SPIR-V,产物是数据不是目标文件 | + +### 4.3 库适配(载荷在 xim,构建面在这里) + +`compat.cublas`、`compat.cudnn`(若许可允许)、`compat.rocblas`、`compat.onemkl` —— +一律 `xpm..deps = { "xim:<载荷>@<上游版本>" }`,自己不探测宿主。 + +--- + +## 5. 设备矩阵 × 验证 lane + +**每一行都有 lane,没有「不做」这一列。** 有硬件用硬件,没硬件用载荷里的模拟器。 + +| 后端 | 编译期判据 | kernel 真跑的 lane | 本机可用? | +|---|---|---|---| +| **CUDA / nvcc** | 无需设备 | **本机 RTX 4080**(compute 8.9) | ✅ 已用过 | +| CUDA(无卡环境) | — | `xim:chipstar` + `xim:pocl` CPU 设备 | ✅ 待接 | +| **SYCL / DPC++ (icpx)** | 无需设备 | ② CUDA 后端 → **本机 4080**,⭐ **已实测跑通**(§6.4.5,`seam: 11 22 33 44`) | ✅ 已验证 | +| **SYCL / AdaptiveCpp** | 无需设备 | `--acpp-targets=omp`,**CPU 上真跑** | ✅ | +| **Vulkan compute** | 无需设备 | ① **lavapipe(已实测跑通)** ② **本机 4080** | ✅ **两条 lane 同机** | +| **OpenCL** | 无需设备 | ① `xim:pocl` CPU ② 本机 `nvidia.icd` | ✅ 两条 | +| **HIP / ROCm** | 无需设备 | ① `HIP_PLATFORM=nvidia` → **本机 4080** ② chipStar+PoCL | ✅ | +| **Metal** | 无需设备 | macOS CI runner | ✅ 四条 macOS CI | +| Level Zero | 无需设备 | 无 Intel GPU ⇒ CPU device | ⚠️ 仅 CPU | + +### 5.1 ⭐⭐ Vulkan 那一行已经跑通了,记在这里作为 lane 的样板 + +本机同时有两个 Vulkan 设备: + +``` +PHYSICAL_DEVICE_TYPE_DISCRETE_GPU NVIDIA GeForce RTX 4080 +PHYSICAL_DEVICE_TYPE_CPU llvmpipe (LLVM 20.1.2, 256 bits) +``` + +lavapipe:Vulkan 1.4、conformance 1.3.1.1、队列含 `QUEUE_COMPUTE_BIT`。 +实跑一次 compute dispatch(`data[i] = data[i]*3+1`,256 元素): + +``` +device: llvmpipe (LLVM 20.1.2, 256 bits) (type=CPU) +checked 256 elements, wrong=0 ; sample: 1 4 7 766 +KERNEL RAN ON CPU AND THE RESULT IS CORRECT +``` + +⭐ **对照也跑了**:只给 `nvidia_icd.json` 时同一程序报 `NO CPU COMPUTE DEVICE` 退 1, +而这台机器确有可用的 4080 ⇒ 对照不是空的,成功确实来自软件设备。 + +⭐⭐ 这就是每条 lane 要达到的样子:**一个真的 dispatch、一个真的结果、一个能让它 +变红的对照。** 不是「设备列出来了」。 + +--- + +## 6. 引擎与插件的分界 + +> **原则:mcpp 核心只做通用构建架构;凡是能由插件回答的,一律由插件回答。** + +「插件」在 mcpp 里已经有名字:**规则包**(`host-module = true` 的包,导出一个只 +`import std` 与 `mcpp` 的模块,发 `mcpp::action` 边)。它已经在 `examples/09-cuda-kernel` +里跑通了。本节的全部工作就是把下面这条线画清楚: + +| 问题 | 归谁 | +|---|---| +| 图里有没有「设备目标」这种节点 | **核心** | +| 一个目标要不要第二次链接、产物是什么角色 | **核心**(边的形状) | +| 那第二次链接的程序叫 `nvlink` 还是 `amdgcn-link` | **插件** | +| 产物身份里有没有 `accel` 这一维、怎么比较 | **核心** | +| `sm_90f` / `gfx942` / `spir64` 怎么拼、怎么算覆盖 | **插件** | +| 一个载荷能不能在这台机器上用 | **核心**(机制) | +| 用什么函数去问驱动要那个数 | **插件** | + +⭐ 判据:**核心里不出现任何厂商名字。** 现有实现已经违反过一次 —— `mcpp.toolchain.devicehost` +读的是 `crt/host_config.h`、`unreachable_device_stage` 调的是 `nvcc`,两者都写死了 +CUDA。⇒ 本轮要把它们改成**由规则包提供探针、核心只负责调用与报告**。 + +### 6.1 核心要补的(通用,不含厂商知识) + +| # | 项 | 通用形式 | +|---|---|---| +| C-1 | **设备目标原语** | `[[target]] kind = "device"`:一种产出「不参与常规链接、由某条边消费」的目标 | +| C-2 | **二次链接边** | 一个目标可声明它需要一次 `role = "device-link"` 的 action,输入是若干对象、输出是一个对象。核心只管顺序与指纹 | +| C-3 | **逐 glob 收窄** | 每条 source glob 可带 `accel` 约束;空集报错、非子集报错 | +| C-4 | **`exclusive` 能力声明** | 包自己声明独占某能力;两个独占同名能力者同时在图里 ⇒ 绑定期报错。**不由核心猜哪些能力互斥** | +| C-5 | **载荷可用性机制** | 载荷描述符可声明「需要宿主某个量 ≥ N」;核心在构建前取那个量并比较,消息给出两边取值与该轴补救。**取值的办法由插件给** | +| C-6 | **含设备代码的归档** | `.a` 携带 `accel` 维并随包传播 | +| C-7 | **`accel` 维的语法是开放的** | 核心只认「后端名 + 版本 + 架构集合 + 可选下界」这个形状,**不内置任何后端的取值表** | + +### 6.2 插件(规则包)要提供的 + +| 规则包 | 提供 | +|---|---| +| `rules-cuda` | nvcc 调用与 `-gencode` 拼装;`sm_XX`/`sm_XXf`/`sm_XXa` 的覆盖判定;`nvlink` 的 device-link 边;`cuDriverGetVersion` 探针;`crt/host_config.h` 的上界读取 | +| `rules-hip` | hipcc 与 `--offload-arch`;`gfx*` 与 generic target 的覆盖判定;`hipRuntimeGetVersion` | +| `rules-sycl` | icpx / acpp 的 `-fsycl-targets` / `--acpp-targets`;形态 B 的整目标编译 | +| `rules-spirv` | `glslc` / `slangc` → SPIR-V;产物是数据不是对象 | +| `rules-metal` | `metal` / `metallib` | + +⚠️ **既有实现的两处要迁走**:`devicehost.cppm` 的 `parse_host_config` 与 doctor 的 +`unreachable_device_stage` 现在都写死 CUDA。改法是核心提供 +「向规则包要一个探针、执行它、报告结果」的通道,**CUDA 的那两个探针搬进 `rules-cuda`**。 +⭐ 收益不是洁癖:AMD/Intel/Apple 各有自己的上界与后端可达性问题,**核心不该长四次**。 + +### 6.3 ⭐ 这条分界对本方案其余部分的影响 + +- §2.1 的驱动配对 ⇒ C-5(机制)+ `rules-cuda`(探针),而不是核心里加一个 CUDA 检查; +- §2.2 的载荷闭包发现 ⇒ 规则包报出缺什么,核心只负责把消息变成一条可执行的建议; +- §7 的九个框架 ⇒ 每一个都只依赖「核心 + 对应规则包」,**不要求核心认识它们**。 + +## 6.3.5 ⭐⭐ 主路线是 LLVM/clang,不是 gcc —— 这是生态事实决定的 + +主次不该按偏好定。逐个核实每个后端的编译器**实体**是什么: + +| 后端 | 编译器 | 是不是 LLVM | +|---|---|---| +| HIP / ROCm | `hipcc` | ✅ **AMD 的 LLVM fork**(`ROCm/llvm-project`) | +| SYCL / DPC++ | `icpx` | ✅ **Intel 的 LLVM fork**(`intel/llvm`,仓库自述 "Home for Intel LLVM-based projects") | +| SYCL / AdaptiveCpp | `acpp` | ✅ 自述 "a powerful, generic **LLVM JIT compiler**" | +| Vulkan / SPIR-V | clspv / slang / glslang 链路 | ✅ | +| Metal | `metal` | ✅ clang 系 | +| OpenMP offload | clang 的更成熟 | ✅ | +| **CUDA** | `nvcc`(EDG 前端) | ❌ **唯一的例外** | + +⭐⭐ 而**唯一的例外也有两条通向 LLVM 的路,两条都已实测**: + +1. nvcc **接受 clang 作 `-ccbin`** —— 本文的编译实验全程用的就是 clang; +2. clang **原生编 CUDA**(`-x cuda`),此时根本没有 nvcc,也就没有宿主编译器上界。 + +### 结论:主次翻转 + +> **加速器路线以 LLVM/clang 为主。gcc 保留支持,但不是默认。** + +理由不是「clang 更好」,是**六个后端里五个本来就是 LLVM,第六个能被 LLVM 编**。 +选 gcc 作主路线意味着为唯一的例外去适配,并在其余五格上一直做转换。 + +具体影响: + +| 处 | 之前 | 改为 | +|---|---|---| +| 加速器工程的默认工具链 | 跟随 mcpp 默认(gcc 16.1.0) | **`llvm@22.1.8` 载荷** | +| CUDA 的默认编译路线 | nvcc + 挑一个上界内的 ccbin | **clang `-x cuda`**(路线 C);nvcc 路线保留给需要 `-gencode` 精细控制或新 CUDA 版本的场合 | +| `rules-*` 的形态 | 每个后端一套厂商驱动 | **五个后端共用 clang 驱动的形状**,差别只在 flag 与 target 拼法 | +| gcc 的位置 | 默认 | **岛形态里设备 TU 的 `-ccbin` 备选**,以及不含设备代码的普通工程照旧 | + +⚠️ **不是弃用 gcc。** mcpp 主体仍以 gcc 自举,普通 C++ 工程不受影响; +改的只是**加速器这条线的默认**。§7.9 场景 1(纯 C++)一个字都不用改。 + +⚠️ 代价要写明:clang 对**新 CUDA 版本**的支持滞后(实测 clang 22 只部分支持到 +CUDA 12.9,clang 18 只到 12.3)。⇒ **需要最新 CUDA 特性的工程走 nvcc 路线**, +由 `rules-cuda` 按「clang 版本 × CUDA 版本」的可用配对来选, +两条路线共存由 xim 的多版本管理承担。 + +--- + +## 6.4 ⭐⭐ 宿主编译器上界不是限制,是四条路里的一次选择(已实测) + +问题:CUDA 13.3 声明 gcc ≤ 15 / clang ≤ 21,CUDA 12.9 声明 gcc ≤ 14, +而 mcpp 的工具链载荷是 **gcc 16.1.0**。这条上界会不会把混合编程钉死? + +**不会。实测了四条路,其中两条把上界整个消掉。** + +### 路线 A —— nvcc 的逃生门 `-allow-unsupported-compiler` + +存在,但它只是关掉检查;真正的风险是新版 libstdc++ 头文件里的构造 EDG 前端解析不了。 +⇒ **可用作最后手段,不作为设计的答案。** + +### 路线 B —— 岛形态把上界关进设备 TU 里(⭐ 已经在用了) + +设备 TU 的 `-ccbin` 与项目的宿主编译器**本来就不必是同一个**: + +``` +项目其余部分 gcc 16.1.0(mcpp 载荷) → .o +设备 TU nvcc -ccbin gcc 15.1.0 → .o ← 上界只约束这一格 + ↑ 也是 xim 载荷 +两者在 extern "C" 的缝上相遇 —— C ABI 跨编译器版本稳定 +``` + +⭐ mcpp 的 registry 里现成就有 **13.3.0 / 15.1.0 / 16.1.0** 三个 gcc 载荷。 +`rules-cuda` 已经在做「挑一个上界内的 ccbin」这件事,只是现在挑的是 host 上的; +改成**从 xim 挑**,这条路就完全在生态内闭合。 + +⚠️ 代价是缝必须干净:`extern "C"`、不跨缝传 std 类型。 +**这正是形态 A 本来就要求的纪律**,不是额外负担。 + +### 路线 C —— 用 clang 直接编 CUDA,上界消失(⭐⭐ 已实测通过) + +clang 原生支持 `-x cuda`。此时**宿主编译器就是 clang 自己**, +「nvcc 接不接受这个宿主编译器」这个问题不存在。 + +实测(本机,2026-09-05): + +``` +$ ~/.xlings/data/xpkgs/xim-x-llvm/22.1.8/bin/clang++ -x cuda t.cu \ + --cuda-path= --cuda-gpu-arch=sm_89 -std=c++17 -c +clang++: warning: CUDA version 12.9 is only partially supported + ✓ 编译成功 +``` + +**全程没有任何宿主编译器上界。** 产物的动态依赖闭包也是干净的: + +``` +libc++.so.1 → xim-x-llvm/22.1.8/lib/... +libc++abi.so.1 → xim-x-llvm/22.1.8/lib/... +libc.so.6 → xim-x-glibc/2.39/lib64/... +``` + +⚠️ **约束是对称的,不是消失的**:nvcc 限制*宿主编译器*版本,clang 限制 *CUDA* 版本。 +实测 clang 18 只到 CUDA 12.3(遇 13.3 报 `fatbinary fatal: Unknown option '-image'`), +clang 22 部分支持到 12.9。⇒ **两条约束都由 xim 的多版本管理承担**, +`rules-cuda` 按「clang 版本 × CUDA 版本」选一对可用的。 + +### 路线 D —— 全部来自载荷,host 只出一个库(已实测) + +上面那个二进制里,glibc、libc++、cudart、nvcc/clang 全部来自载荷; +host 参与的只有 `libcuda.so.1`。 + +⚠️⚠️ **但不能用 `LD_LIBRARY_PATH` 去够它** —— 实测把宿主 `/lib/x86_64-linux-gnu` +放进 `LD_LIBRARY_PATH`,进程 **segfault(exit 139)**,因为宿主 glibc 与载荷 glibc +在同一个地址空间里撞了。**正解就是 `compat.cuda-driver` 现在做的事**: +farm 出**单独一条** `libcuda.so.1` 符号链接,经 `runtime.library_dirs` 进搜索路径。 + +### 路线 D 的补充实测:farm 一条符号链接**不够**,闭包要连它的 `DT_NEEDED` + +把 `libcuda.so.1` farm 进一个目录、写进 RUNPATH,产物闭包确实干净: + +``` +$ ldd tC6 | grep -c '/usr/lib\|/lib/x86_64-linux-gnu' +0 +libc++.so.1 → xim-x-llvm/22.1.8/... libc.so.6 → xim-x-glibc/2.39/... +``` + +⚠️ **但驱动仍然 dlopen 不到。** 宿主的 `libcuda.so.1` 自己声明五个 `DT_NEEDED`: + +``` +libm.so.6 libc.so.6 libdl.so.2 libpthread.so.0 librt.so.1 +``` + +私有 loader 必须能在**同一个闭包里**解析这五个 —— farm 一条链接只解决了 +「找得到 libcuda」,没解决「libcuda 找得到它自己要的东西」。 + +⭐ `compat.cuda-driver` 之所以在真实工程里能用(`examples/09-cuda-kernel` 实测 +输出 `12 24 36 48`),是因为 **mcpp 的链接模型把 glibc 载荷目录也放进了那条搜索 +路径**。手搓 `clang -Wl,-rpath,...` 绕过了这个模型,于是缺了那一半。 + +⚠️⚠️ **由此:判据 C0 必须经 mcpp 跑,不能手搓 clang。** 被测的对象包含 mcpp 的 +链接模型本身;绕过它去测,测的是另一个东西。 +(旁证:手搓产物的 `INTERP` 请求 glibc **2.39** 却解析到 **2.44** —— +一个 mcpp 不会产生的不一致。) + +### 结论 + +| 路线 | 上界还在吗 | 代价 | 定位 | +|---|---|---|---| +| A `-allow-unsupported-compiler` | 关掉检查 | 可能真的编不过 | 最后手段 | +| B 岛形态 + xim 里挑 ccbin | **只约束设备 TU** | 缝要干净(本来就要) | gcc 工程 / 需要 `-gencode` 精细控制时 | +| C clang 编 CUDA | **消失**,换成 CUDA 版本上界 | clang 对新 CUDA 支持滞后 | ⭐ **默认路线**(见 §6.3.5) | +| D 全载荷 | — | 无 | 上面三条的共同底座 | + +⇒ **不是「混合编程受限」,是引擎不该替用户选。** 核心提供「设备 TU 可以有自己的 +宿主编译器」这个通用能力(§6.1 C-1),**由哪个规则包按哪条路线去选,是插件的事。** + +--- + +## 6.4.5 ⭐⭐⭐ 形态 B × modules 已实测:缝形态是**唯一**可行,而它统一了全部后端 + +这是方案里最后一格「不知道会怎样」。已测完(DPC++ 7.1.0 载荷 + 本机 RTX 4080)。 + +### 结果 + +| 形态 | 结果 | +|---|---| +| kernel 写在**模块接口单元**里,SYCL 头进 GMF | ❌ **DPC++ 的 driver 产不出 BMI** | +| **缝形态**:模块接口 SYCL-free,SYCL 在普通 TU | ✅ **跑通**,`seam: 11 22 33 44`,kernel 在 4080 上执行 | + +⭐ 第一行的判据取自**最简可能的模块**(`module;` + `#include ` + +`export module tiny;`),它同样失败: + +``` +clang-offload-bundler: error: 'pcm': invalid file type specified +clang++: error: clang-offload-bundler command failed with exit code 1 +``` + +⇒ **不是 C++ 语义冲突,是 DPC++ 的 driver 管线不认识 `.pcm`。** 这是上游的一个 +洞,不是设计的选择;它会被上游修好,而在那之前形态 B 只能走缝。 + +### ⭐⭐⭐ 由此得到的统一结论 + +> **缝形态不是 CUDA 的权宜,它是全部后端唯一都成立的那一个。** + +- CUDA:已发布并实测(`examples/09-cuda-kernel`,输出 `12 24 36 48`); +- SYCL:**今天唯一能用的形态**(上表); +- HIP / Vulkan / Metal:同构 —— 设备代码进自己的 TU,经 `extern "C"` 与 C++ 相遇。 + +⇒ **一个工程形状覆盖全部后端。** 用户学一次,到处适用; +mcpp 侧也只需要支持一种形状,而不是每个后端一种。 + +### 缝的三条纪律,每条都是实测撞出来的 + +1. **缝必须 `extern "C"`。** 先用 C++ linkage 声明,链接期报 + `undefined reference to 'add_inplace@seam(...)'` —— 模块 linkage 把名字改了。 + ⭐ 与 `examples/09-cuda-kernel` 的纪律**逐字相同**,这不是巧合。 +2. **必须显式指定设备目标。** 不指定 `-fsycl-targets` 时默认编 spir64 镜像, + 而本机只有 CUDA 后端,运行期报 + `No kernel named _ZTSZZ11add_inplace... was found`。 + ⇒ **`accel` 维要记录的正是这个**:产物为哪个 SYCL target 编的。 +3. ⚠️⚠️ **PTX 版本受驱动限制** —— 与 §2.1 同一个失败模式的**第三次独立出现**: + + | SYCL TU 用哪个 CUDA 载荷编 | 运行 | + |---|---| + | CUDA 12.9 | ❌ `CUDA_ERROR_UNSUPPORTED_PTX_VERSION` | + | CUDA 12.4(驱动上限) | ✅ `seam: 11 22 33 44` | + + ⭐ 三次(nvcc 产物、SYCL 产物、载荷选择)都是**编译链接全干净、运行期才说话**。 + 这把 §2.1 从「一条 CUDA 的注意事项」升格为**跨后端的通用约束**: + `accel` 维必须能表达「这个产物要求驱动至少多新」,而不只是「为哪个架构编的」。 + +### ⚠️ 一条前置条件:mcpp 的 LLVM 载荷缺 offload 工具 + +实测 `clang -x cuda -fgpu-rdc`(跨 TU `__device__` 调用)在 `xim:llvm@22.1.8` 上失败: + +``` +clang++: error: unable to execute command: posix_spawn failed: No such file or directory +clang++: error: llvm-offload-binary command failed with exit code 1 +``` + +该载荷是 **slim 构建(36 个二进制)**,不含 `llvm-offload-binary` / +`clang-linker-wrapper` / `clang-offload-bundler`。 +⭐ 对照:**DPC++ 载荷全带**(它就是 LLVM 的一个发行)。 +⇒ **「以 LLVM 为主」的前置条件是给 `xim:llvm` 补上 offload 工具**, +或发一个 `xim:llvm-offload` 伴生包。非 RDC 路径不受影响(已实测通过)。 + +--- + +## 6.4.6 ⭐⭐ CUDA 还是不是「唯一的例外」 + +§6.3.5 的表里 CUDA 是唯一非 LLVM 的一格。综合本节与 §6.4 的实测,答案是: +**在构建层面它已经不是例外了**,理由有三条,每条都已实测: + +1. **clang 原生编 CUDA**(`-x cuda`)—— 此时 CUDA 源码就是 clang 的一种输入语言, + 与 C++/SYCL/HIP 走同一个前端; +2. **SYCL 能编到 CUDA**(`-fsycl-targets=nvptx64-nvidia-cuda`)—— 本节实测, + kernel 真的跑在 4080 上。**LLVM 工具链直达 NVIDIA 硬件**; +3. **chipStar 把 CUDA 编到 SPIR-V**,同样是 LLVM 路径。 + +⇒ **nvcc 从「CUDA 的定义」降级为「CUDA 的一个后端」。** 它仍然必要,但只在三处: + +| 还需要 nvcc 的场合 | 为什么 | +|---|---| +| 最新的 CUDA 版本 | clang 落后(实测 clang 22 部分支持到 12.9) | +| `-gencode` 的精细控制 | `sm_XXf` 家族目标等 clang 未覆盖的拼法 | +| RDC / device link | 在 `xim:llvm` 补齐 offload 工具之前 | + +⭐ **对用户的意义:项目形状、依赖声明、feature、诊断在所有后端上是同一套; +换后端只换 `[features]` 与载荷,不换开发方式。** 这就是「统一的构建与开发体验」的 +具体含义 —— 不是把 nvcc 藏起来,是让它不再决定工程长什么样。 + +⚠️ **兜底也写明**:若某个后端将来出现连缝形态都容纳不了的要求, +它由**自己的规则包**单独支持(§6.2),核心不为它变形。**单独支持是有的, +但它是插件的单独,不是架构的分叉。** + +--- + +## 6.5 ⭐ 多设备后端复用既有通用机制,不新增概念 + +用户体验的关键不是新语法,是**新东西能不能落在已经存在的格子里**。逐项对照: + +| 需求 | 复用哪个既有机制 | 是否新增概念 | +|---|---|---| +| 选后端 | **`[features]`** —— `--features cuda` / `vulkan` / `rocm` | 否 | +| 按后端换依赖 | **`[target.'cfg(accelerator="…")'.dependencies]`** | 否(`accelerator` 已是 layer 键) | +| 按后端换源码 | **`sources` 的逐 glob 约束** | 否(glob 已有,加一个键) | +| 调设备编译器 | **规则包 + `mcpp::action` 边** | 否(`examples/09-cuda-kernel` 已在用) | +| 装工具链 | **xim 载荷 + `xpm..deps`** | 否 | +| 跑/调试/分析 | **`mcpp run` 的 runner 机制** —— `--profile nsys`、`--check sanitizer` | 否 | +| 产物选变体 | **`pack::AbiTag` 的 `accel` 维 + `tag_check`** | 否(已落地) | +| 声明能力独占 | `provides` + 新增 `exclusive` | ⚠️ **一个布尔键** | + +⭐⭐ **整套多设备支持只新增了两样东西**:`accel` 这一个身份维(已落地), +和 `exclusive` 这一个布尔键。其余全部是既有机制的取值扩展。 + +⚠️ 这是本方案对「简洁」的操作性定义:**不是少写字,是不让用户学第二套概念。** +一个已经会用 `--features` 的人,切后端时不需要学任何新东西。 + +### 6.5.1 feature 与 accel 的关系要说清 + +它们**不是同一件事**,混淆会出错: + +- **feature 是意图**:「我要 CUDA 支持」—— 消费者写的,可组合,可传播。 +- **`accel` 是事实**:「这个产物是为 sm_89 编的」—— 构建产生的,进身份,被比较。 + +```toml +[features] +cuda = { accel = "cuda12.9+{sm_80,sm_89}" } # feature 决定 accel,不是等于它 +``` + +⭐ 一个 feature 可以决定 `accel`,但 `accel` 也可以由 `--accel` 直接给, +或由 `[build] accel` 写死。**三个入口一个出口**,与 `--target` / `[toolchain]` +的关系同形。 + +--- + +## 7. 工业级生态覆盖 + +⚠️ **不选「小而完整」的样例包。** 判据是:**工业上真的有人用它出货**, +并且它的构建复杂到能把本方案的每一处都压出来。分五档,每档回答一个不同的问题。 + +### 7.A 推理运行时 —— 「能不能承载真实部署」 + +| 框架 | 版本 | 为什么是它 | 压出什么 | +|---|---|---|---| +| **ONNX Runtime** | 最新 release | 工业推理事实标准;CUDA / TensorRT / oneDNN / OpenVINO / DirectML / CoreML 六个 EP | **多 provider 共存** ⇒ 直接压 §6.4 的 `exclusive`;跨平台 EP ⇒ 压 `cfg(accelerator=)` | +| **llama.cpp / ggml** | llama.cpp `v0.4.0` / ggml `v0.23.0`(均 2026-09-04) | 部署量最大的本地推理栈;**一个代码库里就有 CUDA/Vulkan/SYCL/HIP/Metal/CPU 全部后端** | 整张设备矩阵;变体分发 | +| **libtorch** | 上游 release | C++ 侧工业消费 PyTorch 的实际形态 | **载荷消费**:cuDNN/NCCL/cuBLAS 闭包、私有 loader 下的 RPATH | + +⚠️ 索引里现有的 `ggml-org.llamacpp` pin 在 `b10069`,`sources` 里**只有 `ggml-cpu/*`** —— +CUDA/Vulkan/SYCL/HIP 一个都没有。补齐它就是把矩阵走一遍。 + +⭐ **PyTorch 全源码构建不在第一批**,理由不是难,是**顺序**:它依赖 cuDNN/NCCL/ +oneDNN/oneMKL 全部就位。第一批做 `libtorch` 载荷消费,把闭包与 RPATH 这条路走通; +源码构建排在算子库档之后。 + +### 7.B 算子库与 kernel 编写 —— 「能不能生产算子」 + +| 框架 | 版本 | 压出什么 | +|---|---|---| +| **CUTLASS** | `4.8.0dev`(2026-08-27) | 形态 A 的极限:海量 `.cu` TU、极重模板、**宿主编译器上界**极敏感、`sm_90a` 这类 arch-specific 目标 | +| **oneDNN** | `v3.13.2`(2026-08-26) | PyTorch/TF 底下的那层;CPU + SYCL GPU ⇒ 形态 B 的工业验证 | +| **cuDNN / NCCL / cuBLAS / rocBLAS / oneMKL** | 各自上游 | 纯载荷:不构建,但必须**可被消费**且闭包成立 | + +### 7.C HPC 可移植性 —— 「一份源码能不能跑遍设备」 + +| 框架 | 版本 | 压出什么 | +|---|---|---| +| **Kokkos** | `5.2.1`(2026-08-17) | DOE 级实际用量;CUDA/HIP/SYCL/OpenMP 由构建期选 ⇒ 压 `cfg(accelerator=)` 的成员判定在真实项目里够不够用 | + +### 7.D 视觉与检索 —— 「工业 pipeline 的两端」 + +| 框架 | 版本 | 压出什么 | +|---|---|---| +| **OpenCV** | `5.0.0`(2026-06-06) | 索引已有 `opencv.opencv`;加 CUDA 模块 ⇒ **含设备代码的静态库**(§6.6) | +| **FAISS** | 上游最新 | 工业向量检索;CUDA kernel + 大量模板 | + +### 7.E Vulkan 计算的独立见证 + +| 框架 | 版本 | 压出什么 | +|---|---|---| +| **ncnn** | `20260526` | GPU 路径就是 Vulkan compute,不依赖任何厂商工具链 | + +⭐ 它与 llama.cpp 的 Vulkan 后端**互为独立见证**:两个不相干的工业项目在 +**lavapipe** 上都跑出正确结果,那条无卡 lane 才算真的 —— 单个项目可能碰巧绕开了。 + +### 7.F ⚠️ 这一节的判据不是「编过了」 + +每个框架的判据都必须是**它自己的正确性判据**,不是构建成功: +ONNX Runtime 跑它自己的 model test、llama.cpp 推理出正确 token、CUTLASS 的 GEMM +与参考实现比对、oneDNN 跑它自己的 benchdnn、Kokkos 跑它自己的 unit test、 +OpenCV 跑它自己的 accuracy test、ncnn 的分类结果与 CPU 后端一致。 + +⭐⭐ **上游自带的测试套件是唯一不会被我写偏的判据。** + +--- + +## 7.5 架构横评:与 CMake / xmake / Cargo / Zig / Bazel / Spack / Conan + +⚠️ 本节的立场:**逐条给出可核对的事实,而不是宣称领先。** 落后的地方写在同一张表里。 + +### 7.5.1 核心差异:产物身份里有没有加速器维 + +这是本方案的中心命题,也是最容易被含糊过去的一条。 + +| 系统 | 「为哪个设备架构编的」记在哪里 | 消费时会不会被检查 | +|---|---|---| +| **CMake** | `CMAKE_CUDA_ARCHITECTURES` —— **一个构建变量** | ❌ 产出的 `.a`/`.so` 不携带它;消费方无从得知 | +| **xmake** | `add_cugencodes()` —— 构建配置 | ❌ 同上 | +| **Cargo** | 无。GPU 全在 `build.rs` 里 shell 出去 | ❌ target triple 不编码 GPU 架构;feature 是集合不是兼容维 | +| **Zig** | 无。`-target` 不含 GPU 架构 | ❌ | +| **Bazel** | `--config=cuda` + `cuda_archs` | ❌ 不进 artifact 身份 | +| **Conan** | `settings` 参与 `package_id` —— **机制在,维度不在** | ⚠️ ConanCenter #11448 就是在要这个维度,未合入 | +| **Spack** | `+cuda cuda_arch=90` 是**一等 variant**,进 hash | ✅ **现有系统里最接近的** | +| **Python wheels** | PEP 817/825 Wheel Variants(`namespace::feature::value`) | ⚠️ 正在加,说明这是全行业公认的缺口 | +| **mcpp** | `pack::AbiTag::accel`,与 triple/compiler/stdlib/standard 并列 | ✅ `tag_check` 在选变体时比较,**拒绝时点名维度与两侧取值** | + +⭐ 诚实的结论:**这不是 mcpp 独有的想法** —— Spack 早就有,wheels 正在加。 +mcpp 的差异在于**它是一个编译器驱动的构建系统**,所以那个维度既进包身份, +**也能在编译前就用它拒绝**;Spack 只在求解依赖图时用它,不驱动编译。 + +### 7.5.2 逐维度对比 + +| 维度 | mcpp | CMake | xmake | Cargo | Zig | Bazel | +|---|---|---|---|---|---|---| +| C++20 modules 一等公民 | ✅ 自举即模块 | ⚠️ 2023 起,部分 | ⚠️ 部分 | — | — | ⚠️ 实验 | +| 包管理内置 | ✅ | ❌ 需 vcpkg/conan | ✅ xrepo | ✅ | ⚠️ 新 | ⚠️ 需 rules_* | +| 工具链本身是载荷(hermetic) | ✅ xim 装 gcc/llvm/glibc | ❌ 用 host | ⚠️ 部分 | ⚠️ rustup | ✅ 自带 libc | ✅ 最强 | +| 产物身份含加速器维 | ✅ | ❌ | ❌ | ❌ | ❌ | ❌ | +| 构建程序是真语言 | ✅ `build.mcpp` 是 C++ | ⚠️ CMake DSL | ✅ Lua | ⚠️ `build.rs` | ✅ Zig | ⚠️ Starlark | +| 交叉编译广度 | ⚠️ 追赶中 | ⚠️ 靠工具链文件 | ⚠️ | ⚠️ | ✅ **最强** | ⚠️ | +| 生态广度 | ❌ **远落后** | ✅ **压倒性** | ⚠️ | ✅ | ⚠️ | ⚠️ | +| 大规模远程缓存 | ❌ | ⚠️ | ⚠️ | ⚠️ | ❌ | ✅ **最强** | + +### 7.5.3 ⚠️ mcpp 明确落后的三处,以及本方案对它们做了什么 + +1. **生态广度。** CMake 拥有一切。⇒ 本方案 §7 用**九个工业级框架**换取覆盖面, + 而不是用数量换;判据是这些框架**自己的测试套件**。 +2. **交叉编译广度不如 Zig。** Zig 自带多目标 libc。⇒ mcpp 的路线是 xim 载荷 + (已有 aarch64/musl/mingw/picolibc 多条),方向相同、覆盖面待补,不在本方案范围。 +3. **远程缓存不如 Bazel。** ⇒ 不在本方案范围,但 §6 的 `accel` 维恰好是远程缓存 + key 的必要成分 —— 没有它,跨机器复用缓存会取到为别的架构编的产物。 + +### 7.5.4 语义清晰度:一处可以直接对比的例子 + +同一件事,在 CMake 与 mcpp 里的表达: + +```cmake +# CMake:三个互不相关的机制,各自回答问题的一部分 +set(CMAKE_CUDA_ARCHITECTURES 80;90) # 编给谁 +find_package(CUDAToolkit REQUIRED) # 从 host 找 +target_link_libraries(app CUDA::cublas) # 链什么 +# 「产物为谁而编」不在任何一处被记录,消费者拿到 .a 无从得知 +``` + +```toml +# mcpp:一处声明,进产物身份,消费时被比较 +[build] +accel = "cuda12.9+{sm_80,sm_90f} ptx>=90" + +[dependencies] +cublas = "12.9.x" # 载荷来自 xim,不是 host +``` + +⭐ 差别不在语法长短,在**那句声明去了哪里**:CMake 的进了构建目录, +mcpp 的进了产物身份,并在下一个消费者构建时被 `tag_check` 读出来。 + +## 7.9 场景示例(参考形态,不是最终措辞) + +⚠️ 以下是**参考示例**,用来把前面的抽象落到用户真正会敲的东西上。 +措辞会随实现调整;不变的是每个场景**只需要哪些概念**。 + +### 场景 1 —— 纯 C++,没有设备代码 + +```toml +[package] +name = "app" +version = "0.1.0" +[language] +standard = "c++23" +``` + +⭐ **一个字都不用改,零代价。** `.cu` 不进默认 glob,`accel` 维为空表示不受约束, +`cfg(accelerator=…)` 不匹配任何东西。这是本方案对绝大多数工程的影响面:没有。 + +### 场景 2 —— C++ + 设备代码(形态 A:岛) + +设备代码住在自己的 TU 里,只经 `extern "C"` 头与 C++ 相遇;C++ 侧照常是模块。 + +``` +app/ + src/app.cppm # 模块接口,import 的人只看见这个 + src/kernels/saxpy.cu # 设备 TU,永不被扫描 import,不产 BMI + include/saxpy/saxpy.h # 缝:extern "C",不含任何 CUDA 类型 +``` + +```toml +[package] +name = "app" +version = "0.1.0" +[language] +standard = "c++23" + +[build] +accel = "cuda12.9+{sm_80,sm_89} ptx>=80" + +[build-dependencies] +rules-cuda = { version = "0.1.0", host-module = true } + +[dependencies] +cuda-driver = { namespace = "compat", version = "2026.09.05" } +``` + +```cpp +// build.mcpp —— 构建程序是真的 C++ +import std; +import mcpp; +import rules_cuda; + +int main() { + return rules_cuda::compile(mcpp::sources("src/kernels/*.cu")) ? 0 : 1; +} +``` + +```cpp +// src/app.cppm —— 缝模块:头文件只在这里出现一次 +module; +#include "saxpy/saxpy.h" +export module app; +import std; + +export namespace app { + std::vector scale(std::vector x, float a) { + mcpp_saxpy(static_cast(x.size()), a, x.data()); + return x; + } +} +``` + +⭐ **消费者 `import app;` 即可,永远看不见 CUDA。** 这就是「把设备代码与 C++ 分开 +以支持模块」那个问题的答案:不是不用头文件,是**头文件只出现在缝的内侧一次**。 + +### 场景 3 —— 整目标形态(形态 B:SYCL) + +没有单独的设备 TU;整个目标由 SYCL 编译器编,kernel 就写在 C++ 里。 + +```toml +[toolchain] +compiler = "dpcpp@7.1.0" # 载荷来自 xim,不是 host + +[build] +accel = "sycl:spir64+{cpu,cuda:sm_89}" + +[build-dependencies] +rules-sycl = { version = "0.1.0", host-module = true } +``` + +```cpp +export module vecadd; +import std; +#include // 形态 B 里它在 GMF + +export std::vector add(std::vector a, std::vector b) { + sycl::queue q; + // ... kernel 直接写在这里,同一个 TU +} +``` + +⚠️ **已实测(§6.4.5):kernel 不能写在模块接口单元里** —— DPC++ 的 driver 产不出 +BMI。形态 B 今天只能走**缝形态**:模块接口 SYCL-free,SYCL 代码在普通 TU 里, +经 `extern "C"` 相遇。上面这个写法要等上游修好 `clang-offload-bundler`。 + +### 场景 4 —— 一份源码,多后端(消费者选) + +```toml +[features] +default = ["cpu"] +cpu = [] +cuda = { accel = "cuda12.9+{sm_80,sm_89}" } +vulkan = { accel = "vulkan1.3+{spirv1.6}" } +rocm = { accel = "hip6.4+{gfx1100}" } + +[target.'cfg(accelerator="cuda")'.dependencies] +cublas = "13.5.1" + +[target.'cfg(accelerator="vulkan")'.dependencies] +vulkan = { namespace = "compat", version = "1.4" } +``` + +```console +$ mcpp build --features cuda +$ mcpp build --features vulkan +$ mcpp build --no-accel # 显式要 CPU-only 变体 +``` + +⭐ `cfg(accelerator=…)` **处处是成员判定** ⇒ 同时启用两个后端时 +`all(accelerator="cuda", accelerator="vulkan")` 表示「两者都启用」,不是不可满足。 + +### 场景 5 —— 消费预编译的加速库 + +```toml +[dependencies] +cublas = "13.5.1" # 载荷在 xim,适配在 mcpp-index +llamacpp = { namespace = "ggml-org", version = "b10069", features = ["cuda"] } +``` + +构建时若本地 `accel` 与产物的 `accel` 不相容: + +``` +error: no prebuilt artifact of `llamacpp` matches this build + dimension `accel`: the artifact provides cuda12.9+{sm_90f} + this build requests cuda12.9+{sm_86} + fix: build from source, or `--no-accel` to take the CPU-only variant +``` + +⭐ 消息点名**维度**、**两侧取值**、**该轴可行的补救** —— 三者缺一,用户就得猜。 + +### 场景 6 —— 无卡机器上开发与 CI + +```console +$ xlings install mesa-lavapipe pocl # 模拟器也是载荷 +$ mcpp test --features vulkan # kernel 在 CPU 上真跑 +``` + +```toml +# 在 CI 上把设备固定成软件实现,判据仍是「结果正确」 +[target.'cfg(accelerator="vulkan")'.env] +VK_ICD_FILENAMES = "${xim:mesa-lavapipe}/share/vulkan/icd.d/lvp_icd.json" +``` + +⚠️ **模拟器 lane 的结果必须与硬件 lane 对照过**(§9 C4/C5),否则等于把 +「lavapipe 说对了」当成「它是对的」。 + +### 场景 7 —— 发布多变体 + +```console +$ mcpp pack --accel 'cuda12.9+{sm_80,sm_89}' +$ mcpp pack --accel 'vulkan1.3+{spirv1.6}' +$ mcpp pack --no-accel # CPU-only,必须排第一 +``` + +⚠️⚠️ **CPU-only 变体在描述符里排第一。** 消费方是首次匹配即用的循环, +而不认识 `accel` 键的旧 mcpp 会跳过它 —— 排第一保证它拿到的是能跑的那个。 + +### 场景 8 —— 混合:C++ 模块 + 设备代码 + 第三方算子库 + +这是工业项目的实际形状,三样同时出现: + +```toml +[package] +name = "infer" +version = "0.1.0" +[language] +standard = "c++23" + +[build] +accel = "cuda12.9+{sm_89}" +sources = [ + "src/**/*.cppm", # C++ 模块 + { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }, # 设备 TU,逐 glob 收窄 +] + +[build-dependencies] +rules-cuda = { version = "0.1.0", host-module = true } + +[dependencies] +cublas = "13.5.1" # 厂商算子库 +cutlass = { namespace = "nvidia", version = "4.8.0" } +llamacpp = { namespace = "ggml-org", version = "b10069", features = ["cuda"] } +``` + +```console +$ mcpp build +$ mcpp run --profile nsys # 走 runner,不新增专用命令 +$ mcpp test --check sanitizer # compute-sanitizer 也是 xim 载荷 +``` + +⭐ 注意这里**没有一行提到 host**:nvcc、cudart、cuBLAS、CUTLASS、nsys、 +compute-sanitizer 全部来自 xim;host 参与的只有 `libcuda.so.1`。 + +### 场景 9 —— 诊断 + +```console +$ mcpp self doctor + Checking device toolkit +warning: cuda will refuse this host compiler: gcc 16 exceeds the bound of 15 + stated in /include/crt/host_config.h. +warning: this payload requires driver CUDA >= 13.0; this host reports 12.4. + Device code will build and link, then fail at the first allocation + with "CUDA driver version is insufficient". Install the 12.9 line, or + update the driver. +``` + +⭐⭐ 第二条正是 §2.1 实测到的那次失败 —— **把一个只在运行期才说话的错误, +搬到构建之前**。 + +--- + +## 8. 跨仓库顺序 + +⚠️ 顺序是硬的:**消费者先发布,索引 `latest` 才能动**。 + +``` +⓪ 立即,无依赖 ── 修已发布的错误示范(§0.1) + ⓪a 从全部用户文档删除 `[xlings] deps`(英文 14 处 / 中文 10 处), + `[xlings.workspace]` 是唯一形式 + ⓪b examples/09-cuda-kernel 改用 `[xlings.workspace]` + ⓪c ⚠️ 该示例去 host 化要等 ①a 的载荷 —— 但**文档与字段先改** + +① xim-pkgindex ── 载荷,三批,批内并行 + ①a 编译必需 cuda-nvcc / cudart / nvrtc / nvjitlink / cccl (150 MB) + ①b 设备与模拟器 pocl / mesa-lavapipe / chipstar / spirv-tools / glslang + dpcpp@7.1.0 / adaptivecpp@25.10.0 / hip-runtime@7.14.0 + ①c 算子库与工具 libcublas(+static 分包) / libcufft / cudnn / nccl + cuda-gdb / cuda-sanitizer / nsight-systems / nsight-compute +② mcpp ── 引擎,有序 + 6.1 设备目标 → 6.2 device link → 6.3 glob 收窄 → 6.4 exclusive + → 6.5 驱动配对 → 6.6 含设备代码的静态库 → 6.7 accel 维扩展 +③ mcpp 发布 ── 一个版本,消费者才能用新键 +④ mcpp-index ── 适配面(依赖 ③) + compat.cuda-driver 改名 + repo 改正;compat.vulkan-icd / opencl-icd; + rules-cuda / rules-hip / rules-spirv;compat.cublas / cudnn / nccl / onemkl +⑤ 框架 ── 依赖 ④,按 §7 的档次 + 7.A llamacpp 多后端 → 7.E ncnn → 7.B CUTLASS + oneDNN + → 7.C Kokkos → 7.D OpenCV + FAISS → 7.A ONNX Runtime → 7.A libtorch +``` + +⭐ ① 与 ② **完全并行** —— 载荷不依赖引擎的新键。 +⚠️ ⑤ 内部也有序:llamacpp 先做,因为它一个项目就覆盖矩阵,能最早暴露设计问题。 + +--- + +## 9. 判据表 + +判据取**只有做对了才会出现的读数**;框架档一律用**上游自带的测试套件**。 + +| # | 判据 | lane | +|---|---|---| +| C0 | ⭐ **默认路线**:`xim:llvm@22.1.8` + `xim:cuda-*@12.9.x`,`clang -x cuda` 编出 `sm_89` 并在 4080 上跑出正确结果,**全程零 host 路径**(`ldd` 里不出现 `/usr/lib`) | 本机硬件 | +| C0b | 同一工程切到 nvcc 路线(`rules-cuda` 选 `-ccbin` 为 `xim:gcc@15.1.0`)结果一致 ⇒ **翻转没有废掉 gcc 路线** | 本机硬件 | +| C1 | `xlings install cuda-nvcc@12.9.86` 后 nvcc 可用,且 `mcpp self doctor` **不报** `cicc: not found` —— 全程不碰 host toolkit | 沙箱 | +| C2 | 驱动只到 12.4 的机器上请求 13.x 载荷 ⇒ **构建前拒绝**,消息含两边版本号与该轴补救 | 本机(驱动 550) | +| C3 | `mcpp run --profile nsys` 产出 report;`--check sanitizer` 报出一处故意写的越界 | 本机硬件 | +| C4 | llama.cpp Vulkan 后端在 **lavapipe** 上推理出正确 token | 无卡 | +| C5 | ncnn 在 lavapipe 上分类结果与 CPU 后端一致 | 无卡 | +| C6 | CUTLASS 的 GEMM 与参考实现比对通过 | 本机 4080 | +| C7 | 同一 Kokkos 源码 `--accel cuda…` / `--no-accel` 两份产物**跑它自己的 unit test 都通过**,而 `accel` 不同 | 两条 | +| C8 | 两个声明 `exclusive` 且同名能力的包同时在图里 ⇒ **绑定期报错**,不是链接期重复符号 | 无需设备 | +| C9 | 跨 TU `__device__` 调用链接成功;改 `device-link` 触发重新 prepare | 本机硬件 | +| C10 | llama.cpp 发多变体后,**旧 mcpp 取到 CPU-only 那个**并正常运行 | 无需设备 | +| C11 | oneDNN 的 `benchdnn` 在 DPC++ CPU device 与 CUDA 后端两条 lane 上都通过 | 两条 | +| C18 | ⭐ 缝形态在**全部**后端上是同一个工程形状:CUDA 与 SYCL 两个工程的 `mcpp.toml` 除 `[features]` 与载荷外**逐行相同** | 无需设备 | +| C19 | 不指定设备目标时 mcpp **在构建期**报出「产物没有为任何可用设备编」,而不是留到运行期的 `No kernel named … was found` | 本机硬件 | +| C20 | `accel` 维能表达「要求驱动至少多新」,且 PTX 版本高于驱动时**构建前拒绝**(§6.4.5 纪律 3) | 本机硬件 | +| C12 | ONNX Runtime 同时启用两个 EP ⇒ 由 `exclusive` 决定报错还是共存,**与它自己的文档一致** | 无卡 + 硬件 | +| C13 | OpenCV 含 CUDA 模块的 `.a` 被消费时,`accel` 维随包传播并被 `tag_check` 读出 | 无需设备 | +| C14 | libtorch 载荷消费后,`mcpp` 的运行期闭包校验通过(cuDNN/NCCL/cuBLAS 全部可达,且**不经 host `/usr/lib`**) | 本机硬件 | +| C15 | **核心里 grep 不到任何厂商名字**(`nvcc`/`cuda`/`hip`/`sycl`/`gfx`/`sm_`),它们只出现在规则包里 | 无需设备 | +| C16 | 卸掉 `rules-cuda` 后,`mcpp self doctor` 的设备一节**安静地不报**,而不是报错或崩 | 无需设备 | +| C17 | 只装 `cuda-nvcc` 不装 `libnvvm` ⇒ doctor 报出 `cicc` 并**指名要装哪个 xim 包** | 沙箱 | + +⭐ C2 与 C10 是「答错比不答更坏」的两条:一个放行了跑不了的载荷,一个让旧客户端 +拿到跑不了的产物。**它们要最先有测试。** + +⭐⭐ C4 与 C5 必须**各自在 4080 上再跑一次**:模拟器 lane 的意义在于它与硬件 +lane 结论一致;只跑模拟器,等于把「lavapipe 说对了」当成「它是对的」。 + +--- + +## 10. 风险 + +| # | 风险 | 处置 | +|---|---|---| +| RK-a | ~~`libcublas` 775 MB 过大~~ | **作废。** 已量:shared 是 `libcublas.so` 102.5 MB + `libcublasLt.so` 484.2 MB,static 单独分包。**体积从来不是理由** | +| RK-b | ~~DPC++ release 可能不含 NVIDIA 后端~~ | **作废。** 已核:v7.1.0 Linux 构建配置为 `--cuda --hip`,release notes 测试矩阵含 NVIDIA CUDA BACKEND。Windows 不含 ⇒ 自建,版本仍为 `7.1.0` | +| RK-c | 模拟器与真硬件行为不同(浮点、竞态) | **每条模拟器 lane 必须有一次真硬件对照**(C4/C5) | +| RK-d | lavapipe 慢 | 用最小模型;判据是「结果正确」不是「快」 | +| RK-e | 载荷总体积 | 组件级拆包:普通构建 150 MB,profile 才装 nsight,static 才装 `-static` 包 | +| RK-f | chipStar 覆盖的 CUDA 子集有限 | 它只作**无卡环境的 kernel 行为**见证;`-gencode`/PTX/fatbin 仍走真卡 | +| RK-g | ⚠️ **框架档工作量远大于引擎档** | 顺序上 llamacpp 先行;它若暴露设计问题,后面七个都会受影响 ⇒ 它是 gate 不是第一个任务 | +| RK-h | ⚠️ ONNX Runtime / libtorch 依赖 cuDNN、NCCL,二者许可与 CUDA Toolkit 不同 | 逐个读它们自己的 manifest/EULA,**按 §3.0 的方法由许可字段决定**,不由判断决定 | +| RK-i | ~~形态 B × modules 未实测~~ | **已测,见 §6.4.5。** kernel 进模块接口单元 ❌(DPC++ driver 产不出 BMI);缝形态 ✅(4080 上跑出 `seam: 11 22 33 44`)。⇒ 缝形态是全部后端唯一都成立的形状 | +| RK-l | ⚠️ **`xim:llvm@22.1.8` 是 slim 构建,缺 offload 工具** ⇒ `-fgpu-rdc` 不可用 | 「以 LLVM 为主」的前置条件:补 offload 工具或发 `xim:llvm-offload` 伴生包。非 RDC 路径不受影响 | +| RK-k | ⚠️ **主路线翻转到 clang 的代价是 CUDA 版本滞后** —— 实测 clang 22 部分支持到 12.9、clang 18 只到 12.3 | 两条路线共存,由 `rules-cuda` 按「clang 版本 × CUDA 版本」可用配对选;需要最新 CUDA 特性的走 nvcc。判据 C0b 保证 gcc/nvcc 路线不被废掉 | +| RK-j | ⚠️ **把 CUDA 探针从核心搬进规则包会动已发布的行为** | `mcpp self doctor` 的两条消息是已发布的契约。迁移期核心保留一份回落,规则包在场时以规则包为准;两条路径的消息逐字相同 | + +--- + +## 11. 与前文的关系 + +前文 `2026-09-05-accelerator-support-design.md` 的 §1-§7、§9-§11、§13-§15 仍然成立。 +本文取代的是: + +- **§8.5** 的三档表 —— 第二档补齐(lavapipe 已实测、PoCL/chipStar 已核), + 第三档的「CUDA 只能上真卡」删除,代之以 §5 的矩阵; +- **§12** 的阶段表 —— 「未做」的理由全部作废,代之以 §5 的 lane 与 §9 的判据; +- **§15.8** —— 两条「缺授权」的结论作废,代之以 §4.2 与 §3.1。 diff --git a/.agents/docs/2026-09-05-multi-device-implementation-plan.md b/.agents/docs/2026-09-05-multi-device-implementation-plan.md new file mode 100644 index 00000000..bbf14bb4 --- /dev/null +++ b/.agents/docs/2026-09-05-multi-device-implementation-plan.md @@ -0,0 +1,135 @@ +# 多设备生态:实施计划与任务表 + +对应方案 `2026-09-05-multi-device-ecosystem-design.md`。本文只管**做什么、谁依赖谁、 +判据是什么、现在到哪一步**。方案改了本文跟着改,反之不成立。 + +## 0. 状态总览 + +| 批次 | 仓库 | 状态 | +|---|---|---| +| ⓪ 修已发布的错误示范 | mcpp | 🟡 T0.1/T0.2 ✅,T0.3 待载荷 | +| ① 载荷 | xim-pkgindex | 🟡 **PR #759**(25 个包,已实测) | +| ② 引擎 | mcpp | 🟡 进行中 | +| ③ 发布 | mcpp | ⬜ | +| ④ 适配面 | mcpp-index | ⬜ | +| ⑤ 框架 | mcpp-index | ⬜ | +| ⑥ 生态验证 | 沙箱 | ⬜ | + +图例:⬜ 未开始 / 🟡 进行中 / ✅ 完成并有判据 / ⛔ 阻塞 + +--- + +## 1. 任务表 + +### ⓪ 修已发布的错误示范(mcpp,无依赖,立刻做) + +| # | 任务 | 判据 | 依赖 | +|---|---|---|---| +| T0.1 | 从全部用户文档删除 `[xlings] deps`,`[xlings.workspace]` 为唯一形式 | `git grep -c 'xlings\] deps' docs/` 为 0(中英双份) | — | +| T0.2 | `examples/09-cuda-kernel` 改用 `[xlings.workspace]` | 示例中不出现 `deps =` | T0.1 | +| T0.3 | 该示例去 host 化(`-L/usr/...`、`/usr/local/cuda/bin/nvcc`) | 示例与规则包里 `grep -c '/usr'` 为 0 | T1.1 | + +### ① 载荷(xim-pkgindex,批内并行) + +| # | 任务 | 版本 | 判据 | 依赖 | +|---|---|---|---|---| +| T1.1 | ✅ 24 个 CUDA 组件(编译/运行/调试/分析/算子库) | 12.9 线 + 13.3 线 | ✅ 载荷编 sm_89 并在 4080 上跑出 `12 24 36 48`;两线并存可切换 | — | +| T1.2 | `llvm-offload`(补 slim 载荷缺的 offload 工具) | 22.1.8 | `clang -x cuda -fgpu-rdc` 编链通过 | — | +| T1.3 | ✅ `dpcpp` | 7.1.0 | ✅ `sycl-ls` 报 `[cuda:gpu] NVIDIA CUDA BACKEND` | — | +| T1.4 | `pocl` | 7.2 | `clinfo` 出现 CPU 设备 | — | +| T1.5 | `mesa-lavapipe` | 25.2.8+ | `vulkaninfo` 出现 `PHYSICAL_DEVICE_TYPE_CPU` | — | +| T1.6 | ✅ 随 T1.1 一并落地 | 13.3.x / 2026.x | 静态检查通过;运行判据待 ⑥ | — | +| T1.7 | ✅ `libcublas`+四个算子库(上游归档含 static,未再拆) | 13.5.1.27 | ⚠️ static 分包推迟:上游一个归档同时含 shared 与 static,拆包要重打,先按上游形态发 | — | +| T1.8 | `chipstar` | 1.3.0 | 无卡机器上跑一个 CUDA kernel | T1.4 | +| T1.9 | `adaptivecpp` | 25.10.0 | `--acpp-targets=omp` 跑 kernel | T1.2 | +| T1.10 | `hip-runtime` / `hipcc` | 7.14.0 | `HIP_PLATFORM=nvidia` 跑 kernel | T1.1 | + +### ② 引擎(mcpp,有序) + +| # | 任务 | 判据 | 依赖 | +|---|---|---|---| +| T2.1 | **C-1 设备目标原语** `[[target]] kind = "device"` | 单测 + e2e:设备目标不参与常规链接 | — | +| T2.2 | **C-2 二次链接边** `role = "device-link"` | 跨 TU `__device__` 调用链接成功(C9) | T2.1, T1.2 | +| T2.3 | **C-3 逐 glob 收窄** | 空集/非子集各报错一次 | T2.1 | +| T2.4 | ✅ **C-4 `exclusive` 能力声明** | ✅ e2e 601:独占对被拒并点名双方;**对照** —— 不声明的两个提供者照常共存。3 条单测 + 中英文档 + `exclusive-capability` 进机器接口契约页 | — | +| T2.5 | **C-5 载荷可用性机制** + 探针通道 | 驱动只到 12.4 时请求 13.x ⇒ 构建前拒绝(C2) | — | +| T2.6 | **C-6 含设备代码的归档** | `.a` 的 `accel` 随包传播(C13) | T2.1 | +| T2.7 | **C-7 `accel` 维语法开放** | `vulkan1.3` / `sycl:spir64` / `hip:gfx1100` 可解析比较 | — | +| T2.8 | **把 CUDA 探针搬进规则包** | 核心 grep 不到厂商名字(C15);卸掉规则包 doctor 安静(C16) | T2.5 | +| T2.9 | **`accel` 表达驱动下界** | PTX 版本高于驱动 ⇒ 构建前拒绝(C20) | T2.5, T2.7 | +| T2.10 | **未指定设备目标的构建期诊断** | 报「没有为任何可用设备编」而非运行期(C19) | T2.7 | + +### ③ 发布(mcpp) + +| # | 任务 | 判据 | +|---|---|---| +| T3.1 | 版本号 + CHANGELOG + 发布 | GitHub/GitCode 资产 sha256 一致;`ci-fresh-install` 全绿 | + +### ④ 适配面(mcpp-index,依赖 ③) + +| # | 任务 | 判据 | 依赖 | +|---|---|---|---| +| T4.1 | `compat.cuda-runtime` → `compat.cuda-driver` 改名 + `repo` 改正 | 旧名保留一个跳转期 | T3.1 | +| T4.2 | `compat.vulkan-icd` / `compat.opencl-icd`(缺失时回落载荷) | 无卡机器上 dlopen 到软件实现 | T1.4, T1.5 | +| T4.3 | `rules-cuda` / `rules-hip` / `rules-sycl` / `rules-spirv` 进索引 | 消费者一行依赖即可用 | T3.1, T2.8 | +| T4.4 | `compat.cublas` / `cudnn` / `nccl` / `onemkl` | 闭包校验通过 | T1.7 | + +### ⑤ 框架(mcpp-index,依赖 ④) + +| # | 框架 | 判据(一律用上游自带测试) | 依赖 | +|---|---|---|---| +| T5.1 | llama.cpp 多后端(CPU/CUDA/Vulkan/SYCL/HIP) | 各后端推理出正确 token(C4) | T4.3 | +| T5.2 | ncnn | lavapipe 上分类与 CPU 一致(C5) | T4.2 | +| T5.3 | CUTLASS | GEMM 与参考比对(C6) | T4.3 | +| T5.4 | oneDNN | `benchdnn` 两条 lane(C11) | T1.3 | +| T5.5 | Kokkos | 自带 unit test 两份产物都过(C7) | T4.3 | +| T5.6 | OpenCV CUDA 模块 | accuracy test(C13) | T4.4 | +| T5.7 | FAISS | 自带测试 | T4.4 | +| T5.8 | ONNX Runtime | model test + EP 共存(C12) | T4.4 | +| T5.9 | libtorch 载荷消费 | 闭包校验(C14) | T4.4 | + +### ⑥ 生态验证(沙箱) + +| # | 任务 | 判据 | +|---|---|---| +| T6.1 | 新建 subos,配 CN mirror,装发布物 | 版本自述正确 | +| T6.2 | 逐条跑 C0–C20 | 全绿,且每条模拟器 lane 有硬件对照 | +| T6.3 | 对照:去掉包后判据必须变红 | 每条判据都测到了东西 | + +--- + +## 2. 并行与关键路径 + +``` +⓪ ──────────────────────────────────────────────► (独立,最先完成) + +① 载荷 ────┬─ T1.1 CUDA ──────────┐ + ├─ T1.2 llvm-offload ──┤ + ├─ T1.3 dpcpp ─────────┤ + ├─ T1.4 pocl ──────────┼──► ④ ──► ⑤ ──► ⑥ + ├─ T1.5 lavapipe ──────┤ + └─ T1.6-10 ────────────┘ + │ +② 引擎 ── T2.1 ─ T2.2 ─ … ─ T2.10 ┴─► ③ 发布 ─┘ +``` + +⭐ **关键路径是 ②**(引擎有序,十个任务串行),① 与它完全并行。 +⚠️ ⑤ 的九个框架是工作量主体,T5.1 是 gate —— 它若暴露设计问题,后八个都受影响。 + +--- + +## 3. 动态更新记录 + +方案允许按实际情况改,但每次改必须写下理由。 + +| 日期 | 改了什么 | 理由 | +|---|---|---| +| 2026-09-05 | 建表 | — | +| 2026-09-05 | §3.0 改正:xim **推荐** xlings-res 双镜像,但**允许**第三方 URL template + per-arch sha256 | 读 `xpkg-creater/SKILL.md` §资源选择策略;索引里 247 个配方直连 github。我原先写「xim 不 re-host」过绝对 | +| 2026-09-05 | T0.1/T0.2 完成 | 文档 24 处清零;示例改用 `[xlings.workspace]` | +| 2026-09-05 | T1.1/T1.3/T1.6/T1.7 落地为 xim PR #759(25 包) | 实测通过:载荷编 sm_89 并在 4080 上跑通;两条线并存可切 | +| 2026-09-05 | `cuda-cccl` 用显式 per-version URL 而非模板 | 上游把组件从 `cuda_cccl` 改名为 `cccl`,目录名进 URL,一个模板 404 | +| 2026-09-05 | 配方用 `io.popen` 列文件 | `os.files` 在 `config()` 沙箱里不可用(`attempt to call a nil value`),`llvm.lua` 也用 popen | +| 2026-09-05 | T2.4 完成 | `exclusive` 是列表不是布尔:一个包可提供多项能力而只有部分独占。schema 警告而非报错,因为绑定期那一处才是执行者 | +| 2026-09-05 | ⚠️ 本机 shim 被 #582 剪掉一次 | 25 次 `xlings install` 后 `mcpp` 等 7 个裸名 shim 消失,store 完好。重装即恢复 —— 又一次受控复现,补进 issue | +| 2026-09-05 | `libcublas` 暂不拆 static | 上游一个归档同时含 shared 与 static,拆分需要重打包并 re-host,与「不 re-host」冲突;先按上游形态发,拆分单列 | diff --git a/docs/05-mcpp-toml.md b/docs/05-mcpp-toml.md index 407bfb06..0876f2da 100644 --- a/docs/05-mcpp-toml.md +++ b/docs/05-mcpp-toml.md @@ -1318,7 +1318,8 @@ runner = ["qemu-aarch64-static"] The rules, for `mcpp run` and `mcpp test` alike: - **A declared runner is used.** Its first element is located by mcpp: first in - the `bin/` directory of each payload declared under `[xlings] deps` (§2.13), + the `bin/` directory of each payload declared under `[xlings.workspace]` + (§2.13), then on `PATH`. A bare name on `PATH` resolves to an xvm shim, which answers for the current SubOS rather than for the package; the payload lookup is what lets a runner name a program the project declared. @@ -1334,14 +1335,14 @@ The rules, for `mcpp run` and `mcpp test` alike: the manifest has no axis to carry; a project whose runner was written for x86_64 developers is still readable on an aarch64 machine. -Provisioning the emulator through `[xlings] deps` is the form for a CI job or -a project built on one host class. `qemu-user-aarch64` in the index is built -for x86_64 Linux only, and `[xlings] deps` provisions on every host that builds -the project, so the entry is written per platform (§2.13): +Provisioning the emulator through `[xlings.workspace]` is the form for a CI job +or a project built on one host class. `qemu-user-aarch64` in the index is built +for x86_64 Linux only, and the table provisions on every host that builds the +project, so the entry is written per platform (§2.13): ```toml -[xlings] -deps = [{ linux = "qemu-user-aarch64" }] +[xlings.workspace] +"xim:qemu-user-aarch64" = { linux = "" } # present on Linux, any version [target.aarch64-linux-musl] runner = ["qemu-aarch64-static"] @@ -1533,6 +1534,46 @@ ambiguity to pin: pinning replaces an error that names both candidates with one that names a mangled symbol. Interchangeable *libraries* (BLAS implementations, which export distinct symbol sets and are selected per link) are unaffected. +#### `exclusive` — a package declaring it is the only provider + +The paragraph above describes a defect the engine cannot detect. Seeing that two +providers define the same symbols requires their object files, which do not +exist when capabilities are bound; and refusing every duplicate provider as a +rule would break the BLAS case in the same paragraph, which is legitimate. + +So the package says it: + +```toml +[package] +name = "compat.cublas" +provides = ["gpu-blas"] +exclusive = ["gpu-blas"] +``` + +Two packages that both provide `gpu-blas`, where at least one declares it +exclusive, are refused when capabilities are bound — before anything is +compiled, naming the capability and both providers: + +``` +error: capability 'gpu-blas' is provided by more than one package, and they + declare it EXCLUSIVE. + providers: [compat.cublas, compat.rocblas] + exclusive: [compat.cublas, compat.rocblas] + Two implementations of one interface define the same symbols, so the + link would resolve every call to whichever archive it reached first. + Keep one of them — a `[capabilities]` pin selects a provider for a + REQUIREMENT and cannot make two definitions of one symbol safe. +``` + +The refusal reports `exclusive-capability` in `--format json` (chapter 11). + +**It is a claim about this package's own symbols**, so an entry that names a +capability the package does not provide is reported as a schema warning: there +is nothing to be exclusive about. And a capability nobody declares exclusive +behaves exactly as before — two BLAS implementations still coexist, and the +existing "two or more, unpinned" error still applies only when something +*requires* the capability. + ### 2.8.2 `[feature-deps.]` — dependencies a feature pulls in A dependency declared under `[feature-deps.]` is **optional**: it is @@ -2526,8 +2567,8 @@ build_finished = "mcpp-hooks-audioplayer niulai-mm" build_failed = "mcpp-hooks-audioplayer niulai-niulai" side_effect = false -[xlings] -deps = ["xim:mcpp-hooks-audioplayer@0.0.1"] +[xlings.workspace] +"xim:mcpp-hooks-audioplayer" = "0.0.1" ``` A different sound for a successful or failed build. `side_effect = false` is diff --git a/docs/07-build-mcpp.md b/docs/07-build-mcpp.md index eca7eef4..34d6c91c 100644 --- a/docs/07-build-mcpp.md +++ b/docs/07-build-mcpp.md @@ -104,7 +104,7 @@ int main() { | `mcpp::dep_bin(pkg, tool)` *(2026.8.5.1+)* | reads `MCPP_DEP__BIN_` — the absolute path of a **host tool** built by a dependency (see below) | | `mcpp::link_script(p)` *(2026.8.19+)* | `mcpp:link-script=` | | `mcpp::runner(tok)` *(2026.8.19.2+)* | `mcpp:runner=` — see below | -| `mcpp::xpkg_dir(ns, name)` / `mcpp::xpkg_dir(name)` *(2026.8.19+)* | the payload directory of a package this manifest declared in `[xlings] deps`; `""` when it was not declared or is not installed (see below) | +| `mcpp::xpkg_dir(ns, name)` / `mcpp::xpkg_dir(name)` *(2026.8.19+)* | the payload directory of a package this manifest declared in `[xlings.workspace]`; `""` when it was not declared or is not installed (see below) | | `mcpp::warning(text)` *(2026.8.21.2+)* | `mcpp:warning=` — see below | | `mcpp::action{…}.submit()` *(2026.8.5.1+)* | `mcpp:action=` — declares a **build-graph node** instead of doing the work here (see below) | @@ -199,7 +199,7 @@ package whose content is implementation-neutral. Asking follows whatever the compiler payload or through the runtime binding, and nothing has to look for it. -### Finding an `[xlings] deps` payload: `xpkg_dir` (2026.8.19+) +### Finding an `[xlings.workspace]` payload: `xpkg_dir` (2026.8.19+) `dep_dir` answers for **mcpp** dependencies. An xlings package is a different namespace with a different store layout, and `xpkg_dir` is the interface for it: diff --git a/docs/11-machine-output.md b/docs/11-machine-output.md index 47566b9a..39f8eae0 100644 --- a/docs/11-machine-output.md +++ b/docs/11-machine-output.md @@ -377,6 +377,7 @@ a program classifying the outcome reads `reason`: | `os-mismatch` | the requested and resolved triples name different systems | | `layer-requirement` | a package requires a layer the resolution did not give it | | `layer-ordering` | the five layers do not stack | +| `exclusive-capability` | two packages provide one capability and at least one declared it exclusive | | `other` | a refusal whose branch has not been given a token yet | ⚠️ **Exit 0 whenever the question was answered, including "refused".** "Would diff --git a/docs/13-baremetal.md b/docs/13-baremetal.md index f33d6cd6..cadc744a 100644 --- a/docs/13-baremetal.md +++ b/docs/13-baremetal.md @@ -291,9 +291,9 @@ fact, selection is a board fact.** | Board-support package | Which startup object and libraries to select, which linker script, which emulator invocation | `-lcrt0-semihost`, `picolibcpp.ld`, `qemu-system-riscv64 -machine virt …` | The middle row is what keeps a package from having to name a C library. -Earlier versions of both ecosystem packages declared -`[xlings] deps = ["xim:picolibc-riscv@1.8.12"]`, which bound a package to one -libc, one architecture and one compiler implementation. That declaration is no +Earlier versions of both ecosystem packages named a libc package in the +environment table directly, which bound a package to one libc, one architecture +and one compiler implementation. That declaration is no longer required, and the target's sysroot column replaced it. A second board on the same ISA is a change of the three values in the bottom @@ -641,8 +641,8 @@ it needs under `[xlings.workspace]`, exports one C++ module for consumers, and emits its board facts from `build.mcpp`. **A declaration there provisions the package on the first build** (since -2026.8.29; the table is `[xlings.workspace]` since 2026.9.3.1, and the older -`[xlings] deps` still works and says so). It is also what lets `mcpp::xpkg_dir` answer *"where +2026.8.29; the table is `[xlings.workspace]`). It is also what lets +`mcpp::xpkg_dir` answer *"where did that package land"*. Both halves matter: the same declaration installs the emulator and tells the build program where it went. diff --git a/docs/15-openkal-cross.md b/docs/15-openkal-cross.md index 29577e50..7ca05a8a 100644 --- a/docs/15-openkal-cross.md +++ b/docs/15-openkal-cross.md @@ -197,8 +197,8 @@ package that provides the emulator is declared for the hosts that can install it: ```toml -[xlings] -deps = [{ linux = "qemu-user-aarch64" }] +[xlings.workspace] +"xim:qemu-user-aarch64" = { linux = "" } [target.aarch64-linux-musl] runner = ["qemu-aarch64-static"] diff --git a/docs/17-the-project-environment.md b/docs/17-the-project-environment.md index 3a516004..de8b9365 100644 --- a/docs/17-the-project-environment.md +++ b/docs/17-the-project-environment.md @@ -8,7 +8,9 @@ and in CI, whatever else those two machines happen to have installed. ```toml [xlings] subos = "tools" -deps = ["xim:qemu-riscv@9.2.4-1"] + +[xlings.workspace] +"xim:qemu-riscv" = "9.2.4-1" ``` Working project: `examples/07-project-subos/`. @@ -85,7 +87,7 @@ should beat a file. ## 3. What the declaration does not decide -`[xlings] deps` names packages to be present in the environment, and each one's +`[xlings.workspace]` names packages to be present in the environment, and each one's payload directory is delivered separately as `MCPP_XPKG__DIR`. That is a different question from `PATH` and stays a different answer: a build program that needs a package's data files (protoc's well-known `.proto` files, say) @@ -132,7 +134,7 @@ gives the full rule. Against that: an isolated environment is a directory that has to be created and populated, and the first build pays for it. Since 2026.8.29 mcpp does that -work — a declared `[xlings] deps` is provisioned on first use, and a named +work — a declared `[xlings.workspace]` entry is provisioned on first use, and a named `[xlings] subos` that does not exist yet is created rather than refused — but the cost is real: the first build on a clean machine downloads and installs before it compiles anything. A project whose tools are ordinary and whose diff --git a/docs/18-devices.md b/docs/18-devices.md index a42f828b..362499d7 100644 --- a/docs/18-devices.md +++ b/docs/18-devices.md @@ -59,7 +59,7 @@ mcpp::run_exclusive(); // this target's runs cannot overlap ``` ⭐ **Name the program, not its path.** mcpp locates it: the `bin/` of a payload -declared under `[xlings] deps` by **any package in the graph** — the consuming +declared under `[xlings.workspace]` by **any package in the graph** — the consuming project first, then its dependencies — and then `PATH`. A board-support package is precisely the thing that knows which emulator or probe reaches its machine, so it declares that payload itself and the consumer declares nothing. Writing an diff --git a/docs/20-accelerators.md b/docs/20-accelerators.md index 614af3a3..bab3b090 100644 --- a/docs/20-accelerators.md +++ b/docs/20-accelerators.md @@ -173,6 +173,31 @@ Family targets and portable forms are what keep the variant matrix finite. Publishing one artifact per chip does not scale; publishing one per generation does. +### The grammar is open + +`cuda`, `hip`, `vulkan` and `sycl` are not a closed set. A backend name, a +version, an architecture list and an optional floor are the whole shape, and +mcpp compares them without a table of who exists: + +``` +vulkan1.3+{spirv1.6} floor>=1.4 +sycl2020+{spir64,nvptx64-sm_89} +hip6.4+{gfx942} +``` + +An architecture whose spelling carries no leading number — `gfx942` — is +compared by equality, because there is no ordering to read out of it. That is +the answer rather than a guess: reading `942` out of the middle would invent a +level AMD does not define. + +`floor>=` is the backend-neutral spelling of the portable-form floor. `ptx>=` +is CUDA's word for the same field and remains accepted, so descriptors written +before this are unaffected; a backend whose portable form is SPIR-V writes +`floor>=` instead of borrowing NVIDIA's term. + +Which spellings mean what for a given backend is the business of that backend's +rule package. What the engine holds is the shape and the comparison. + When nothing matches, the refusal names the dimension and both sides: ``` diff --git a/docs/zh/05-mcpp-toml.md b/docs/zh/05-mcpp-toml.md index 47265dee..c5745284 100644 --- a/docs/zh/05-mcpp-toml.md +++ b/docs/zh/05-mcpp-toml.md @@ -1151,7 +1151,7 @@ runner = ["qemu-aarch64-static"] 规则对 `mcpp run` 与 `mcpp test` 相同: -- **声明了 runner 就使用它。** 其第一个元素由 mcpp 定位:先在 `[xlings] deps`(§2.13) +- **声明了 runner 就使用它。** 其第一个元素由 mcpp 定位:先在 `[xlings.workspace]`(§2.13) 声明的每个载荷的 `bin/` 目录里找,再找 `PATH`。`PATH` 上的裸名会命中 xvm shim,而 shim 按当前 SubOS 而非按包作答;先查载荷,runner 才能直接写工程声明过的程序名。 - **声明的 runner 找不到或启动不了是错误**,错误里带程序名、搜索过的目录和 errno。 @@ -1163,13 +1163,13 @@ runner = ["qemu-aarch64-static"] 这个三元组在本机是原生的 —— 清单没有承载它的轴;为 x86_64 开发者写的 runner 在 aarch64 机器上仍可用。 -通过 `[xlings] deps` 装模拟器是 CI 任务或单一宿主类别工程的形态。索引里的 -`qemu-user-aarch64` 只为 x86_64 Linux 构建,而 `[xlings] deps` 在每台构建本工程的 -宿主上都会 provisioning,所以条目按平台写(§2.13): +通过 `[xlings.workspace]` 装模拟器是 CI 任务或单一宿主类别工程的形态。索引里的 +`qemu-user-aarch64` 只为 x86_64 Linux 构建,而这张表在每台构建本工程的宿主上都会 +provisioning,所以条目按平台写(§2.13): ```toml -[xlings] -deps = [{ linux = "qemu-user-aarch64" }] +[xlings.workspace] +"xim:qemu-user-aarch64" = { linux = "" } # Linux 上存在即可,版本不限 [target.aarch64-linux-musl] runner = ["qemu-aarch64-static"] @@ -1320,6 +1320,37 @@ ld: obj/mcpplibs_pa/src/impl.o: in function `cap_probe': 缺陷**而非可 pin 的歧义:pin 会把一个点名两个候选的报错,换成一个点名 mangled 符号的报错。 可互换的**库**(各 BLAS 实现导出不同的符号集合,按链接选其一)不受此影响。 +#### `exclusive` —— 包声明自己是唯一提供者 + +上一段描述的是一个引擎**看不见**的缺陷:要看出两个 provider 定义了同一批符号, +需要它们的目标文件,而绑定 capability 时那些还不存在;而「一律拒绝重复 provider」 +又会打断同一段里那个合法的 BLAS 用例。 + +所以由包自己声明: + +```toml +[package] +name = "compat.cublas" +provides = ["gpu-blas"] +exclusive = ["gpu-blas"] +``` + +两个都提供 `gpu-blas` 的包,只要其中至少一个声明了独占,就在绑定 capability 时 +被拒绝 —— 在任何东西被编译之前,并点名该能力与双方: + +``` +error: capability 'gpu-blas' is provided by more than one package, and they + declare it EXCLUSIVE. + providers: [compat.cublas, compat.rocblas] + exclusive: [compat.cublas, compat.rocblas] +``` + +该拒绝在 `--format json` 里报 `exclusive-capability`(见第 11 章)。 + +**它是关于这个包自己的符号的声明**,所以一条指向本包并不提供的能力的条目会被报为 +schema 警告:那里没有可独占的东西。而无人声明独占的能力行为完全不变 —— 两个 BLAS +实现照常共存,既有的「两个或更多、未 pin」报错也仍然只在**有人 require** 该能力时出现。 + ### 2.8.2 `[feature-deps.]` —— 由 feature 拉取的依赖 在 `[feature-deps.]` 下声明的依赖是**可选的**:仅当该 feature 激活时(根 `--features`, diff --git a/docs/zh/07-build-mcpp.md b/docs/zh/07-build-mcpp.md index 222644d9..17dc8482 100644 --- a/docs/zh/07-build-mcpp.md +++ b/docs/zh/07-build-mcpp.md @@ -97,7 +97,7 @@ int main() { | `mcpp::dep_bin(pkg, tool)` *(2026.8.5.1+)* | 读 `MCPP_DEP__BIN_` —— 依赖构建出的 **host 工具**的绝对路径(见下) | | `mcpp::link_script(p)` *(2026.8.19+)* | `mcpp:link-script=` | | `mcpp::runner(tok)` *(2026.8.19.2+)* | `mcpp:runner=` —— 见下 | -| `mcpp::xpkg_dir(ns, name)` / `mcpp::xpkg_dir(name)` *(2026.8.19+)* | 本 manifest 在 `[xlings] deps` 里声明的包的载荷目录;没声明或没安装时返回 `""`(见下) | +| `mcpp::xpkg_dir(ns, name)` / `mcpp::xpkg_dir(name)` *(2026.8.19+)* | 本 manifest 在 `[xlings.workspace]` 里声明的包的载荷目录;没声明或没安装时返回 `""`(见下) | | `mcpp::warning(text)` *(2026.8.21.2+)* | `mcpp:warning=` —— 见下 | | `mcpp::action{…}.submit()` *(2026.8.5.1+)* | `mcpp:action=` —— **声明一个构建图节点**,而不是在这里把活干了(见下) | @@ -177,7 +177,7 @@ const char* sr = mcpp::sysroot_dir(); // 目标的 C 库根目录,没有则 宿主目标上 `sysroot_dir()` 为空:那里 C 库随编译器载荷或运行时绑定而来,没人需要找它。 -### 找到 `[xlings] deps` 的载荷:`xpkg_dir`(2026.8.19+) +### 找到 `[xlings.workspace]` 的载荷:`xpkg_dir`(2026.8.19+) `dep_dir` 回答的是 **mcpp** 依赖。xlings 包是另一个命名空间、另一套 store 布局, `xpkg_dir` 是它的接口: diff --git a/docs/zh/11-machine-output.md b/docs/zh/11-machine-output.md index 5ba42ad1..d39f0b49 100644 --- a/docs/zh/11-machine-output.md +++ b/docs/zh/11-machine-output.md @@ -331,6 +331,7 @@ mcpp why toolchain [--target ] [--toolchain ] --format json | `os-mismatch` | 请求的与解析出的三元组指向不同的操作系统 | | `layer-requirement` | 某个包要求的层,解析没有给出 | | `layer-ordering` | 五层叠不起来 | +| `exclusive-capability` | 一个能力有多个提供者,而其中至少一个声明了独占 | | `other` | 一处还没有被命名的拒绝分支 | ⚠️ **只要问题被回答了就退 0,包括答案是「拒绝」。** 「它能不能构建,不能的话 diff --git a/docs/zh/13-baremetal.md b/docs/zh/13-baremetal.md index 5d0970f8..a705a41a 100644 --- a/docs/zh/13-baremetal.md +++ b/docs/zh/13-baremetal.md @@ -252,9 +252,8 @@ extern "C" int main() { | 目标 | 用哪个编译器与哪份 C 库,两者都从目标的表行解析并按需安装 | `pin = llvm@22.1.8`、`sysroot = xim:picolibc-riscv@1.8.12` | | 板级支持包 | 选哪个启动对象与哪些库、哪份链接脚本、哪条模拟器命令行 | `-lcrt0-semihost`、`picolibcpp.ld`、`qemu-system-riscv64 -machine virt …` | -中间那一行正是让包不必指名 C 库的原因。两个生态包早先都声明过 -`[xlings] deps = ["xim:picolibc-riscv@1.8.12"]`,这把包绑死在一份 libc、一种架构 -与一种编译器实现上。该声明已不再需要,目标的 sysroot 列取代了它。 +中间那一行正是让包不必指名 C 库的原因。两个生态包早先都在环境表里直接指名过一份 +libc 包,这把包绑死在一份 libc、一种架构与一种编译器实现上。该声明已不再需要,目标的 sysroot 列取代了它。 同一 ISA 上的第二块板子,是把最后一行里的三个取值换掉。它不需要引擎作任何改动。 @@ -567,8 +566,8 @@ error: no runner is configured for 'riscv64-none-elf' — a freestanding artifac 板级支持包是一个普通的 mcpp 包。它在 `[xlings.workspace]` 下声明所需的模拟器,为消费者 导出一个 C++ 模块,并从 `build.mcpp` 发出它的板级事实。 -**那里的声明会在首次构建时供给该包**(2026.8.29 起;这张表自 2026.9.3.1 起是 -`[xlings.workspace]`,旧的 `[xlings] deps` 仍然生效并会说明这一点)。它同时让 +**那里的声明会在首次构建时供给该包**(2026.8.29 起;这张表是 +`[xlings.workspace]`)。它同时让 `mcpp::xpkg_dir` 能回答「那个包落在哪」。两半都要紧:同一条声明既装上模拟器,也告诉 构建程序它装到了哪。 diff --git a/docs/zh/15-openkal-cross.md b/docs/zh/15-openkal-cross.md index a0c2fa5d..7bb08e67 100644 --- a/docs/zh/15-openkal-cross.md +++ b/docs/zh/15-openkal-cross.md @@ -169,8 +169,8 @@ hosted 交叉目标使用同一个键,配用户态模拟器(2026.9.2.1)。在 x8 按能安装它的宿主声明: ```toml -[xlings] -deps = [{ linux = "qemu-user-aarch64" }] +[xlings.workspace] +"xim:qemu-user-aarch64" = { linux = "" } [target.aarch64-linux-musl] runner = ["qemu-aarch64-static"] diff --git a/docs/zh/17-the-project-environment.md b/docs/zh/17-the-project-environment.md index cea034c7..46c042a8 100644 --- a/docs/zh/17-the-project-environment.md +++ b/docs/zh/17-the-project-environment.md @@ -7,7 +7,9 @@ ```toml [xlings] subos = "tools" -deps = ["xim:qemu-riscv@9.2.4-1"] + +[xlings.workspace] +"xim:qemu-riscv" = "9.2.4-1" ``` 可运行的工程:`examples/07-project-subos/`。 @@ -72,7 +74,7 @@ SubOS 里。前置让被声明的环境成为默认答案;其余的仍在它后 ## 3. 这条声明不决定什么 -`[xlings] deps` 声明的是「环境里要有哪些包」,而每个包的载荷目录另有通道交付, +`[xlings.workspace]` 声明的是「环境里要有哪些包」,而每个包的载荷目录另有通道交付, 即 `MCPP_XPKG__DIR`。这与 `PATH` 是两个问题,答案也保持分开:需要某个包 的数据文件(比如 protoc 自带的 well-known `.proto`)的构建程序问目录,需要 **运行**某个程序的构建程序问 `PATH`。 @@ -108,7 +110,7 @@ create/bootstrap that environment instead of falling back to active/default 败;私有环境让这个问题不成立。 代价一侧:隔离环境是一个必须被创建并填充的目录,而这笔账由首次构建来付。 -2026.8.29 起 mcpp 会做这件事 —— 声明在 `[xlings] deps` 里的包在首次使用时被供给, +2026.8.29 起 mcpp 会做这件事 —— 声明在 `[xlings.workspace]` 里的包在首次使用时被供给, 一个尚不存在的具名 `[xlings] subos` 会被创建而不是被拒绝 —— 但代价是实打实的: 干净机器上的第一次构建会先下载安装,然后才编译。工具很普通、版本也无所谓的项目, 不声明、直接继承机器的那份更划算。 diff --git a/docs/zh/18-devices.md b/docs/zh/18-devices.md index b495bebc..714f9d7e 100644 --- a/docs/zh/18-devices.md +++ b/docs/zh/18-devices.md @@ -48,7 +48,7 @@ mcpp::run_exclusive(); // 这个目标的运行不能重叠 ``` ⭐ **写程序名,不要写路径。** mcpp 会定位它:先找**图中任何一个包**在 -`[xlings] deps` 里声明的载荷的 `bin/`(消费工程优先,然后是它的依赖),再找 `PATH`。 +`[xlings.workspace]` 里声明的载荷的 `bin/`(消费工程优先,然后是它的依赖),再找 `PATH`。 板级包正是那个知道「哪个模拟器或探针能抵达这台机器」的东西,所以由它声明, **消费者什么都不用声明**。用 `mcpp::xpkg_dir` 拼绝对路径是多余的,而且引入了一个 失败模式 —— **声明不是安装**,查询可能返回空,于是没有配置任何 runner 而没有任何 diff --git a/docs/zh/20-accelerators.md b/docs/zh/20-accelerators.md index e5e29444..a056129c 100644 --- a/docs/zh/20-accelerators.md +++ b/docs/zh/20-accelerators.md @@ -163,6 +163,26 @@ error: mcpplibs.gpuonly@0.1.0: no prebuilt artifact matches this toolchain. 这正是该维度存在所要搬移的那个失败。没有它,构建干净地链接完成, 而程序在第一次 kernel 启动时失败,消息里既没有包名,也没有任何一侧期望的架构。 +### 语法是开放的 + +`cuda`、`hip`、`vulkan`、`sycl` 不是一个封闭集合。**后端名 + 版本 + 架构集合 + +可选下界**就是全部形状,mcpp 在没有「谁存在」这张表的情况下比较它们: + +``` +vulkan1.3+{spirv1.6} floor>=1.4 +sycl2020+{spir64,nvptx64-sm_89} +hip6.4+{gfx942} +``` + +拼法里不带前导数字的架构 —— 如 `gfx942` —— 按**相等**比较,因为那里读不出任何序。 +这是答案而不是回避:从中间抠出 `942` 会发明一个 AMD 并未定义的等级。 + +`floor>=` 是可移植形式下界的**后端中立**拼法。`ptx>=` 是 CUDA 对同一字段的说法, +仍然被接受,所以此前写下的描述符不受影响;可移植形式是 SPIR-V 的后端写 `floor>=`, +不必借用 NVIDIA 的词。 + +各拼法对某个后端具体意味着什么,是**那个后端的规则包**的事。引擎持有的是形状与比较。 + ### 发布多个变体 一个包可以发布多个产物,消费者取第一个标签接受它的。**把 CPU-only 的产物排在第一位。** diff --git a/examples/09-cuda-kernel/app/mcpp.toml b/examples/09-cuda-kernel/app/mcpp.toml index fd718e10..8ca590cf 100644 --- a/examples/09-cuda-kernel/app/mcpp.toml +++ b/examples/09-cuda-kernel/app/mcpp.toml @@ -29,8 +29,8 @@ cuda-runtime = "2026.09.05" # mcpp a path it can put on the artifact's runtime search path — mcpp's private # loader does not consult /usr/lib, so without this the statically linked CUDA # runtime cannot dlopen the driver and reports it as missing. -[xlings] -deps = [{ linux = "libcuda-host-link" }] +[xlings.workspace] +"xim:libcuda-host-link" = { linux = "0.0.1" } [build] # The CUDA runtime is linked STATICALLY. mcpp refuses a dynamic link against diff --git a/modules/manifest/src/toml.cppm b/modules/manifest/src/toml.cppm index 19596a6f..fc4279f3 100644 --- a/modules/manifest/src/toml.cppm +++ b/modules/manifest/src/toml.cppm @@ -815,6 +815,29 @@ std::expected parse_string(std::string_view content, m.unknownCapabilities.push_back(entry); m.requires_ = *v; } + // [package] exclusive — capabilities this package claims sole provision of. + // + // Not validated against the reserved prefix: exclusivity is a property of + // the package's own symbols, and the capabilities it applies to are the + // feature system's open vocabulary. What IS checked is that each entry is + // something this package actually provides — an entry naming a capability + // the package does not supply cannot be enforced and is almost certainly a + // typo, so it is reported rather than ignored. + if (auto v = doc->get_string_array("package.exclusive")) { + for (auto const& entry : *v) { + bool supplied = std::ranges::find(m.provides, entry) != m.provides.end(); + for (auto const& [_, provs] : m.featureProvides) + if (!supplied) + supplied = std::ranges::find(provs, entry) != provs.end(); + if (!supplied) + m.schemaWarnings.push_back(std::format( + "[package] exclusive names '{}', which this package does not " + "provide. Exclusivity is a claim about this package's own " + "symbols, so it can only apply to a capability listed in " + "`provides` (or in a feature's `provides`).", entry)); + } + m.exclusive = *v; + } // std-module / std-compat-module / std-module-flags. // // ⚠️ THEY BELONG UNDER `[build]`, AND `[package]` IS THE OLDER SPELLING. diff --git a/modules/manifest/src/types.cppm b/modules/manifest/src/types.cppm index 896964ff..519f832d 100644 --- a/modules/manifest/src/types.cppm +++ b/modules/manifest/src/types.cppm @@ -1347,6 +1347,37 @@ struct Manifest { // through untouched, exactly as they do in `provides`. // ⚠️ The spelling is `requires_` because `requires` is a keyword. std::vector requires_; + // [package] exclusive — the capabilities this package claims it is the ONLY + // provider of. + // + // WHY THIS IS DECLARED RATHER THAN INFERRED + // + // Two packages providing one capability is usually fine and sometimes the + // point: `compat.openblas` and an MKL package both provide `blas`, and a + // build that links one program against each is legitimate. So the engine + // cannot refuse duplicate providers as a rule — the rule would break a case + // this project documents. + // + // What it also cannot do is detect the case that is NOT fine. Two + // implementations of one accelerator interface define the same symbols, and + // the link then resolves every call to whichever archive the linker reached + // first. Seeing that requires the object files, which do not exist when + // capabilities are bound. + // + // So the package says it. An author who knows their library defines + // `cublasCreate` writes: + // + // provides = ["gpu-blas"] + // exclusive = ["gpu-blas"] + // + // and two such packages in one graph are refused at binding time, naming + // both — instead of at link time naming a symbol, or at run time naming + // nothing at all. + // + // ⚠️ An entry that is not also in `provides` (or a feature's `provides`) is + // a typo and is reported: claiming exclusivity over something you do not + // provide cannot be acted on. + std::vector exclusive; // [package] std-module / std-module-flags — a package that IS a standard // library says where its `std' module source is and what that source needs // to compile. The build tool otherwise asks the COMPILER where std.cppm is diff --git a/src/build/prepare.cppm b/src/build/prepare.cppm index 4dfdedca..2416aa45 100644 --- a/src/build/prepare.cppm +++ b/src/build/prepare.cppm @@ -6473,6 +6473,10 @@ prepare_build(bool print_fingerprint, // apply() as each package's features activate; bound after the loops below. std::map> capProviders; std::vector> capRequires; + // Who claimed sole provision of what. Separate from capProviders because + // the question it answers is different: capProviders asks "can this + // requirement be satisfied", this asks "can these two coexist at all". + std::map> capExclusive; { auto sanitize = [](std::string f) { for (auto& c : f) @@ -6494,6 +6498,7 @@ prepare_build(bool print_fingerprint, // active. Requirements are bound after all packages are processed. const auto& pcap = pkg.manifest.package.name; for (auto& cap : pkg.manifest.provides) capProviders[cap].push_back(pcap); + for (auto& cap : pkg.manifest.exclusive) capExclusive[cap].push_back(pcap); for (auto& f : active) { if (auto it = pkg.manifest.featureProvides.find(f); it != pkg.manifest.featureProvides.end()) @@ -7484,6 +7489,41 @@ prepare_build(bool print_fingerprint, p = c + 1; } + // EXCLUSIVE CAPABILITIES, CHECKED BEFORE REQUIREMENTS ARE BOUND. + // + // Ordering is deliberate. A requirement conflict is reported by naming + // the requirement; this one exists whether or not anything requires the + // capability, because the defect is that two implementations of one + // interface are in the same link. Reporting it first means the message + // names the real problem rather than a symptom of it. + for (auto const& [cap, claimers] : capExclusive) { + auto it = capProviders.find(cap); + if (it == capProviders.end()) continue; + std::vector providers; + for (auto const& p : it->second) + if (std::find(providers.begin(), providers.end(), p) == providers.end()) + providers.push_back(p); + if (providers.size() < 2) continue; + + std::string list, claimed; + for (auto const& p : providers) list += (list.empty() ? "" : ", ") + p; + for (auto const& c : claimers) claimed += (claimed.empty() ? "" : ", ") + c; + refusal::record(refusal::Code::ExclusiveCapability); + return std::unexpected(std::format( + "capability '{}' is provided by more than one package, and {} " + "declares it EXCLUSIVE.\n" + " providers: [{}]\n" + " exclusive: [{}]\n" + " Two implementations of one interface define the same " + "symbols, so the link would\n" + " resolve every call to whichever archive it reached " + "first. Keep one of them —\n" + " a `[capabilities]` pin selects a provider for a " + "REQUIREMENT and cannot make two\n" + " definitions of one symbol safe.", + cap, claimers.size() == 1 ? "it" : "they", list, claimed)); + } + std::set boundCaps; for (auto& [cap, requirer] : capRequires) { if (!boundCaps.insert(cap).second) continue; // one diagnosis per cap diff --git a/src/build/refusal.cppm b/src/build/refusal.cppm index 73901ed3..74a15594 100644 --- a/src/build/refusal.cppm +++ b/src/build/refusal.cppm @@ -59,6 +59,11 @@ enum class Code { LldRequiredAbsent, // the row links through lld directly and none is here HostToolToolchain, // build.mcpp under a cross target has no host toolchain StdModulePrecompile, // the std module could not be precompiled + // Two packages provide one capability and at least one of them declared it + // exclusive. Distinct from CapabilityPin: that one is about SELECTING a + // provider for a requirement, this one is about two implementations of one + // interface being in the same link at all. + ExclusiveCapability, Other, // a refusal that has not been given a code yet }; @@ -79,6 +84,7 @@ constexpr std::string_view name(Code c) { case Code::LldRequiredAbsent: return "lld-required-absent"; case Code::HostToolToolchain: return "host-tool-toolchain"; case Code::StdModulePrecompile: return "std-module-precompile"; + case Code::ExclusiveCapability: return "exclusive-capability"; case Code::Other: return "other"; } return "other"; diff --git a/src/pack/abi_tag.cppm b/src/pack/abi_tag.cppm index 447ce516..8fcd8da8 100644 --- a/src/pack/abi_tag.cppm +++ b/src/pack/abi_tag.cppm @@ -68,6 +68,11 @@ struct AccelSet { // embeds PTX so later hardware can JIT; AMD has no equivalent and obtains // the same reach through family targets on the `archs` side instead. An // empty floor must therefore widen nothing. + // The floor of the artifact's embedded PORTABLE form, below which it no + // longer reaches. Named for CUDA's PTX because that is where the idea and + // the published spelling come from; the field itself is backend-neutral and + // `floor>=` is its neutral spelling on the wire. A backend with no portable + // form leaves it empty, which widens nothing -- AMD is that case. std::string ptxFloor; }; @@ -287,8 +292,16 @@ std::vector parse_accel(std::string_view s) { j = k + 1; } } - if (auto pf = tail.find("ptx>="); pf != std::string_view::npos) - a.ptxFloor = std::string(trim_sv(tail.substr(pf + 5))); + // The floor below which an artifact's embedded portable form no + // longer reaches. `ptx>=` is CUDA's spelling of it and the one + // already published; `floor>=` is the backend-neutral one, so a + // backend whose portable form is not PTX -- SPIR-V, say -- does not + // have to borrow NVIDIA's word for it. One field, two spellings, + // and the value means the same thing to the comparison either way. + if (auto pf = tail.find("floor>="); pf != std::string_view::npos) + a.ptxFloor = std::string(trim_sv(tail.substr(pf + 7))); + else if (auto pf2 = tail.find("ptx>="); pf2 != std::string_view::npos) + a.ptxFloor = std::string(trim_sv(tail.substr(pf2 + 5))); } out.push_back(std::move(a)); } diff --git a/tests/e2e/601_exclusive_capability.sh b/tests/e2e/601_exclusive_capability.sh new file mode 100755 index 00000000..c8bd800f --- /dev/null +++ b/tests/e2e/601_exclusive_capability.sh @@ -0,0 +1,81 @@ +#!/usr/bin/env bash +# An exclusive capability with two providers is refused at binding time. +# +# Two packages providing one capability is usually fine and sometimes the +# point: an OpenBLAS package and an MKL package both provide `blas`, and a +# build that links one program against each is legitimate. What is not fine is +# two implementations of ONE accelerator interface in one link: they define the +# same symbols, and the link resolves every call to whichever archive it +# reached first. +# +# The engine cannot tell those apart -- seeing the symbol overlap needs object +# files that do not exist when capabilities are bound -- so the package says it. +# This asserts both halves: the declared-exclusive pair is refused, and the pair +# that declares nothing still builds. +set -e + +TMP=$(mktemp -d) +trap "rm -rf $TMP" EXIT +cd "$TMP" + +mkpkg() { # $1 name $2 capability $3 "exclusive"|"" + mkdir -p "$1/src" + { + echo '[package]' + echo "name = \"$1\"" + echo 'version = "0.1.0"' + echo "provides = [\"$2\"]" + [ -n "$3" ] && echo "exclusive = [\"$2\"]" + echo '[language]' + echo 'standard = "c++23"' + echo '[targets.'"$1"']' + echo 'kind = "lib"' + } > "$1/mcpp.toml" + echo "int ${1//-/_}_anchor(void) { return 0; }" > "$1/src/$1.c" +} + +consumer() { # $@ dep names + rm -rf app; "$MCPP" new app > /dev/null; cd app + { + echo '[package]' + echo 'name = "app"' + echo 'version = "0.1.0"' + echo '[language]' + echo 'standard = "c++23"' + echo '[dependencies]' + for d in "$@"; do echo "$d = { path = \"../$d\" }"; done + } > mcpp.toml + rm -f src/*.cppm + echo 'int main(){return 0;}' > src/main.cpp +} + +# ── Half one: two exclusive providers of one capability are refused ───────── +mkpkg gpublas-a gpu-blas exclusive +mkpkg gpublas-b gpu-blas exclusive +consumer gpublas-a gpublas-b + +if "$MCPP" build > out.log 2>&1; then + cat out.log + echo "FAIL: two exclusive providers of 'gpu-blas' were accepted" + exit 1 +fi +grep -q "gpu-blas" out.log || { cat out.log; echo "FAIL: refusal does not name the capability"; exit 1; } +grep -q "gpublas-a" out.log || { cat out.log; echo "FAIL: refusal does not name the first provider"; exit 1; } +grep -q "gpublas-b" out.log || { cat out.log; echo "FAIL: refusal does not name the second provider"; exit 1; } +grep -qi "exclusive" out.log || { cat out.log; echo "FAIL: refusal does not say why"; exit 1; } +echo "PASS: exclusive pair refused, naming the capability and both providers" + +# ── Half two: the SAME graph without the claim still builds ──────────────── +# +# The control matters. Without it this test would also pass if mcpp refused +# every duplicate provider, which is the behaviour the `exclusive` key exists +# to avoid: it would break the documented OpenBLAS/MKL case. +cd "$TMP" +mkpkg blas-a blas +mkpkg blas-b blas +consumer blas-a blas-b + +"$MCPP" build > ok.log 2>&1 || { cat ok.log; echo "FAIL: two ordinary providers were refused"; exit 1; } +echo "PASS: two ordinary providers of one capability coexist" + +echo "PASS: exclusive capability" diff --git a/tests/unit/test_abi_tag_accel.cpp b/tests/unit/test_abi_tag_accel.cpp index a0648670..d82d2518 100644 --- a/tests/unit/test_abi_tag_accel.cpp +++ b/tests/unit/test_abi_tag_accel.cpp @@ -159,3 +159,63 @@ TEST(AbiTagAccel, UnparseableTextMeansNoDeviceCode) { EXPECT_TRUE(mcpp::pack::parse_accel("(none)").empty()); EXPECT_TRUE(mcpp::pack::parse_accel("12.8+{sm_90}").empty()); // no backend } + +// ── The grammar is open: backends mcpp has never heard of ────────────────── +// +// The `accel` dimension exists to carry "which device was this built for", and +// which devices exist is not mcpp's question. These assert that a spelling from +// a backend the engine knows nothing about round-trips and compares -- and, +// just as importantly, that an unknown arch shape falls back to equality rather +// than to some numeric guess. + +TEST(AccelOpenGrammar, VulkanSpirvRoundTrips) { + auto a = mcpp::pack::parse_accel("vulkan1.3+{spirv1.6}"); + ASSERT_EQ(a.size(), 1u); + EXPECT_EQ(a[0].backend, "vulkan"); + EXPECT_EQ(a[0].version, "1.3"); + ASSERT_EQ(a[0].archs.size(), 1u); + EXPECT_EQ(a[0].archs[0], "spirv1.6"); + EXPECT_EQ(mcpp::pack::accel_str(a), "vulkan1.3+{spirv1.6}"); +} + +TEST(AccelOpenGrammar, SyclTargetRoundTrips) { + auto a = mcpp::pack::parse_accel("sycl2020+{spir64,nvptx64-sm_89}"); + ASSERT_EQ(a.size(), 1u); + EXPECT_EQ(a[0].backend, "sycl"); + ASSERT_EQ(a[0].archs.size(), 2u); + EXPECT_EQ(a[0].archs[0], "spir64"); + EXPECT_EQ(a[0].archs[1], "nvptx64-sm_89"); +} + +TEST(AccelOpenGrammar, AnUnknownArchShapeComparesByEquality) { + // `gfx942` carries no leading number, so there is no ordering to read out + // of it. Equality is the only question available and it is the one asked -- + // rather than reading `942` out of the middle and inventing a level. + auto pub = mcpp::pack::parse_accel("hip6.4+{gfx942}"); + auto same = mcpp::pack::parse_accel("hip6.4+{gfx942}"); + auto diff = mcpp::pack::parse_accel("hip6.4+{gfx1100}"); + EXPECT_TRUE (mcpp::pack::accel_accepts(pub, same)); + EXPECT_FALSE(mcpp::pack::accel_accepts(pub, diff)); +} + +TEST(AccelOpenGrammar, FloorHasABackendNeutralSpelling) { + // `ptx>=` is CUDA's word and the one already published. `floor>=` says the + // same thing without borrowing it, so a backend whose portable form is + // SPIR-V does not have to write `ptx`. + auto neutral = mcpp::pack::parse_accel("vulkan1.3+{spirv1.6} floor>=1.4"); + ASSERT_EQ(neutral.size(), 1u); + EXPECT_EQ(neutral[0].ptxFloor, "1.4"); + + auto cuda = mcpp::pack::parse_accel("cuda12.9+{sm_80} ptx>=80"); + ASSERT_EQ(cuda.size(), 1u); + EXPECT_EQ(cuda[0].ptxFloor, "80"); +} + +TEST(AccelOpenGrammar, TheExistingPtxSpellingStillParses) { + // The negative control for the change above: adding a second spelling must + // not have moved the first. Descriptors published before it exist. + auto a = mcpp::pack::parse_accel("cuda12.8+{sm_80,sm_90f} ptx>=90"); + ASSERT_EQ(a.size(), 1u); + EXPECT_EQ(a[0].ptxFloor, "90"); + EXPECT_EQ(mcpp::pack::accel_str(a), "cuda12.8+{sm_80,sm_90f} ptx>=90"); +} diff --git a/tests/unit/test_manifest.cpp b/tests/unit/test_manifest.cpp index ad9e36f2..965a1b0e 100644 --- a/tests/unit/test_manifest.cpp +++ b/tests/unit/test_manifest.cpp @@ -4853,3 +4853,65 @@ hardware = {} if (w.find("hardwear") != std::string::npos) sawTypo = true; EXPECT_TRUE(sawTypo) << "no warning for [feature-xlings.hardwear]"; } + +// ── [package] exclusive ──────────────────────────────────────────────────── +// +// The key is a claim about THIS package's own symbols, so it can only apply to +// a capability the package supplies. The parse therefore accepts the list and +// reports an entry that is not supplied — the value is still carried, because +// the binding-time check is the one that acts on it and a warning must not +// silently change what was declared. + +TEST(ManifestExclusive, CarriesTheListAndWarnsAboutNothing) { + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "compat.cublas" +version = "0.1.0" +provides = ["gpu-blas", "gpu-fft"] +exclusive = ["gpu-blas"] +)"); + ASSERT_TRUE(m.has_value()) << m.error().format(); + ASSERT_EQ(m->exclusive.size(), 1u); + EXPECT_EQ(m->exclusive[0], "gpu-blas"); + std::string all; + for (auto const& w : m->schemaWarnings) { all += w; all += '\n'; } + EXPECT_TRUE(m->schemaWarnings.empty()) << all; +} + +TEST(ManifestExclusive, ClaimingSomethingNotProvidedIsReported) { + // Not an error: the declaration is carried so the binding-time check sees + // exactly what was written. What it cannot do is act on a capability this + // package does not supply, and a silent no-op is how a typo survives. + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "compat.cublas" +version = "0.1.0" +provides = ["gpu-blas"] +exclusive = ["gpu-bals"] +)"); + ASSERT_TRUE(m.has_value()) << m.error().format(); + ASSERT_EQ(m->exclusive.size(), 1u); + std::string all; + for (auto const& w : m->schemaWarnings) { all += w; all += '\n'; } + EXPECT_NE(all.find("gpu-bals"), std::string::npos) + << "the warning must name the entry that cannot be enforced:\n" << all; +} + +TEST(ManifestExclusive, AFeatureProvidedCapabilityCounts) { + // `provides` is not the only way a package supplies a capability, and a + // check that only looked there would warn about a correct manifest. + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "compat.cublas" +version = "0.1.0" +exclusive = ["gpu-blas"] + +[features] +cuda = { provides = ["gpu-blas"] } +)"); + ASSERT_TRUE(m.has_value()) << m.error().format(); + std::string all; + for (auto const& w : m->schemaWarnings) { all += w; all += '\n'; } + EXPECT_EQ(all.find("gpu-blas"), std::string::npos) + << "a feature-provided capability is supplied; no warning is due:\n" << all; +}