diff --git a/.agents/docs/2026-09-05-multi-device-ecosystem-design.md b/.agents/docs/2026-09-05-multi-device-ecosystem-design.md index 8ca935d20..19dbdd1c1 100644 --- a/.agents/docs/2026-09-05-multi-device-ecosystem-design.md +++ b/.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 时自己就定义它。 diff --git a/.agents/docs/2026-09-05-multi-device-implementation-plan.md b/.agents/docs/2026-09-05-multi-device-implementation-plan.md index bbf14bb4e..67ca15d18 100644 --- a/.agents/docs/2026-09-05-multi-device-implementation-plan.md +++ b/.agents/docs/2026-09-05-multi-device-implementation-plan.md @@ -7,12 +7,12 @@ | 批次 | 仓库 | 状态 | |---|---|---| -| ⓪ 修已发布的错误示范 | mcpp | 🟡 T0.1/T0.2 ✅,T0.3 待载荷 | -| ① 载荷 | xim-pkgindex | 🟡 **PR #759**(25 个包,已实测) | -| ② 引擎 | mcpp | 🟡 进行中 | +| ⓪ 修已发布的错误示范 | mcpp | ✅ T0.1/T0.2 合入 `1e2137b`;T0.3 在 **#563** | +| ① 载荷 | xim-pkgindex | ✅ **#759 已合入 `a4644a7`**,15 项 CI 全绿,25 个包;⚠️ 13.x 后端不可达由 **#760** 修 | +| ② 引擎 | mcpp | ✅ T2.1–T2.10 全部完成并各有判据 | | ③ 发布 | mcpp | ⬜ | -| ④ 适配面 | mcpp-index | ⬜ | -| ⑤ 框架 | mcpp-index | ⬜ | +| ④ 适配面 | mcpp-index | 🟡 T4.1/T4.4 ✅(**#347 已合入 `8a9ca64`**);T4.2 待 ① 的 pocl/lavapipe;T4.3 待 ③ | +| ⑤ 框架 | mcpp-index | 🟡 T5.1 走到「载荷版本不匹配」:见下方记录。它作为 gate 已经交付了它该交付的东西 —— 暴露出 C-6 的引擎缺口 | | ⑥ 生态验证 | 沙箱 | ⬜ | 图例:⬜ 未开始 / 🟡 进行中 / ✅ 完成并有判据 / ⛔ 阻塞 @@ -27,14 +27,14 @@ |---|---|---|---| | 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 | +| T0.3 | ✅ 示例去 host 化 | ✅ `mcpp build -v` 的命令行里 `/usr` CUDA 路径 **0 处**;nvcc 与两处 include 全来自 `xpkgs/xim-x-cuda-*`;`mcpp run` → `12 24 36 48` | 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.2 | ⬜ `llvm-offload` | 22.1.8 | ⚠️ 已有替代:`dpcpp@7.1.0` 载荷自带全套 offload 工具,需要 RDC 的工程可用它。独立 `llvm-offload` 仍待做 | — | | 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` | — | @@ -48,16 +48,16 @@ | # | 任务 | 判据 | 依赖 | |---|---|---|---| -| 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.1 | ✅ **C-1 设备目标原语** —— 以既有 `mcpp::action` 的 `role = "artifact"` 落地,不新增 target kind | ✅ artifact 角色的产物不进链接(ninja_backend 既有);理由见「动态更新记录」 | — | +| T2.2 | ✅ **C-2 二次链接边** —— 以「object 角色的 action 以其它 action 的产物为输入」落地,核心只管顺序与指纹 | ✅ e2e 607:artifact 角色产物被 object 角色消费,中间产物不进链接线,顺序由图保证。⚠️ C9(nvcc `-rdc=true` + `-dlink` 真机)本机做不到 —— 见下方记录 | T2.1 | +| T2.3 | ✅ **C-3 逐 glob 收窄** `sources = [{ glob, accel }]` | ✅ e2e 606 四段:覆盖 ⇒ 编译且设备源到达构建程序;`--no-accel` ⇒ 整条 glob 排除;非子集 ⇒ 拒绝并点名两侧(`accel-mismatch`);空集 ⇒ 拒绝点名 glob。6 条单测 | — | | 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 | +| T2.5 | ✅ **C-5 载荷可用性机制**:探针通道 `mcpp::fact` / `mcpp::floor`(协议 v7),核心只比较;根工程的构建程序说完后再查一次 | ✅ e2e 605:根 build.mcpp 陈述的下界被比较并拒绝(两侧取值 + `version-floor-unmet`);对照:满足则构建 | — | +| T2.6 | ✅ **C-6 含设备代码的归档** | ✅ e2e 608:object 角色 action 的产物进 `.a`(判据是 `ar t` 的成员表与 `nm` 的符号,不是退出码 —— 空档案也会成功)。缺陷由 T5.1 暴露:llama.cpp 的 305 个 `.cu` 全被丢弃并只留一条警告 | T2.1 | +| T2.7 | ✅ **C-7 `accel` 维语法开放**(#562) | ✅ 5 条单测 `AccelOpenGrammar.*`;`floor>=` 为中性拼法 | — | +| T2.8 | ✅ **把 CUDA 探针搬进规则包** | ✅ `test_core_vendor_probes`:剥注释后 `src/` 无厂商工具名(自带分母,枚举 < 100 文件即判失败);doctor 的设备节与 `mcpp.toolchain.devicehost` 一并删除;同样的读数由 rules-cuda 产出 | T2.5 | +| T2.9 | ✅ **`accel` 表达驱动下界** | ✅ 实测:`fact=cuda.driver=12.4` + `floor=cuda.driver >= 13.0` ⇒ 13.3 工具包在编译前被拒;PTX 高于驱动 ⇒ 警告(点名架构集合仍可运行) | T2.5, T2.7 | +| T2.10 | ✅ **未指定设备目标的构建期诊断** | ✅ `accel` 不含架构 ⇒ 规则包在提交任何 action 前拒绝并说明;`--no-accel` 走 CPU 实现而非「编了但没设备码」 | T2.7 | ### ③ 发布(mcpp) @@ -69,10 +69,10 @@ | # | 任务 | 判据 | 依赖 | |---|---|---|---| -| T4.1 | `compat.cuda-runtime` → `compat.cuda-driver` 改名 + `repo` 改正 | 旧名保留一个跳转期 | T3.1 | +| T4.1 | ✅ `compat.cuda-runtime` → `compat.cuda-driver` 改名 + `repo` 改正 | ✅ 旧条目冻结保留;工作区成员 `tests/examples/cuda-driver` **同时依赖新旧两个名字**,让跳转期这条承诺有判据(此前它只是一句注释) | — | | 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 | +| T4.4 | 🟡 `compat.cudart` + `cublas`/`cufft`/`curand`/`cusolver`/`cusparse` ✅ | ✅ e2e 判据=新成员 `tests/examples/cuda-curand`:无卡机器断言库能加载并应答,有卡再断言 [0,1] 与均值。⚠️ `cudnn`/`nccl`/`onemkl` 仍缺 xim 载荷 | T1.7 | ### ⑤ 框架(mcpp-index,依赖 ④) @@ -130,6 +130,32 @@ | 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.5 的驱动取数**撤出核心** | 仓库自带的 `test_runtime_contract` 不变量禁止「厂商词 + 探针启动」在 `src/` 共现,抓住了我写的 `nvidia-smi` 调用。**这条规则先于本工作存在且是对的** —— 关系留下并单测,取数改由声明抵达,归规则包通道 | +| 2026-09-05 | T0.3 完成,并暴露一处「几乎为真」 | nvcc 自动加的是**它自己的** `../include`,12.x 那里有 `crt/` 却没有 `cuda_runtime.h`(在 cudart 组件)。不显式传载荷 include ⇒ nvcc 从 `/usr/include` 取头、连带读宿主的 `host_config.h`,**用着载荷的编译器却报宿主工具包的错** | +| 2026-09-05 | ① 全部完成并合入 | 25 个载荷;五个命令(nvcc/cuda-gdb/ncu/nsys/sycl-ls)从 shim 可达 | | 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」冲突;先按上游形态发,拆分单列 | +| 2026-09-05 | ⚠️⚠️ **C9(RDC 真机)在本机不可测,判据改为通用形状** | nvcc 路线在这台机器上**两个工具包都用不了**:12.9 满足驱动而不满足 C 库,13.3 满足 C 库而不满足驱动。C9 需要 nvcc 才能构造,于是它的判据退回到 e2e 607 所测的**通用链式 action**(与厂商无关的同一形状),真机 RDC 留给有 13.x 驱动的机器 | +| 2026-09-05 | ⚠️⚠️ **12.9 工具包 + glibc ≥ 2.41 是一个不能配的对** | 12.9 的 `crt/math_functions.h` 为宿主重声明 C23 的 `cospi`/`sinpi`/`rsqrt` **不带 `noexcept`**,glibc 2.41+ 带,而 C++17 起它是函数类型的一部分 ⇒ 六条 `exception specification is incompatible`,点名两个头文件而不给结论。13.3 不再重声明,同一份 C 库下干净通过。规则包读 `bits/mathcalls.h` 直接拒这一对 | +| 2026-09-05 | **引擎新增两个环境变量而不是复用 `MCPP_TARGET_SYSROOT`** | 后者是**档位**事实(裸机 C 库载荷,宿主目标为空),已有消费者按它判断档位;第二个编译器要的是**环境**事实。复用会把两个含义压进一个键 —— 正是 [[a-value-acquires-the-receiving-layers-requirements]] 那条 | +| 2026-09-05 | **`-B` 的守卫收敛为 `gcc::binutils_prefix_dir`,并只对 GCC 作答** | 三处副本(registry / gcc / flags),其中一处注释写着「Mirrors the guard in build/flags.cppm」。clang 的命令行本就不带 `-B`,所以对 clang 作答会描述一个没人传的开关 | +| 2026-09-05 | ⚠️ **构建程序 helper 在 Linux 上改用 `DT_RPATH`** | RUNPATH 只对 helper **自己**的 needed 生效;它 `dlopen` 的宿主库的依赖(`libdl.so.2`)按私有 loader 的默认搜索,搜不到。这是规则包能读到驱动版本的前提 | +| 2026-09-05 | ⚠️⚠️ **① 的 13.x 载荷装完不能用** —— xim #760 | nvcc 用 `$(TOP)/nvvm/bin/cicc` 找自己的后端,而 13.x 把 `nvvm/` 与 `crt/` 拆成了独立包=独立载荷根。**载荷完整、`nvcc --version` 正常、组件都装了**,编译时 `exit 127`。修法在 `install()` 里把这两个目录链回来;`os.exists`/`os.ln` 在配方沙箱里都不存在,`os.cp(symlink=true)` 是「保留源里的符号链接」而不是「建一个」 | +| 2026-09-05 | ⚠️⚠️ **「没有加速器」被写成了显示用的 `(none)`** | `accel_str` 为空集打印 `(none)` 是给 ABI 标签读的;`resolvedAccel` 把这个拼法当值传了出去 ⇒ ①`MCPP_ACCEL=(none)` 到达**每一个从未提过加速器的工程**,与手册承诺的空串矛盾;②指纹里 `if (!accel.empty())` 恒真,给所有工程都追加了 `#accel=(none)`。判据只能靠构建程序**写文件**取得 —— 它的 stdout 只在非零退出时才打印。e2e 605 第四段标题写着「变量与 layer 都清空」却只测了 layer,这就是它逃过套件的原因 | +| 2026-09-05 | ④ 的 T4.1/T4.4 落地为 mcpp-index #347 并合入 | 六个新包 + 一次改名;两处上游耦合写进配方(`crt/` 在编译器组件里;NVIDIA 的 `.so` 带 `RUNPATH=$ORIGIN` 会关掉继承的 RPATH ⇒ 要一并 farm glibc 三个存根) | +| 2026-09-05 | ⚠️ **path 索引里包的命名空间由「索引名」决定,而不是描述符里的 `namespace`** | 本地验证时 `[indices] localidx = { path = ... }` 下 `compat.cudart` 解析不到,而诊断说「a package with this name exists under another namespace: compat.cudart」—— 把索引名改成 `compat` 即通。诊断本身值得单独修 | +| 2026-09-05 | ⭐⭐ **T5.1 作为 gate 立刻兑现了**:它暴露了 C-6 的引擎缺口 | llama.cpp 的 CUDA 后端是 305 个 `.cu` 挂在 `kind = "lib"` 上。引擎把 object 角色的 action **只**挂到可执行/共享库/测试上,于是每一个 action 都被丢弃、只留一条警告,而构建**成功**并产出一个不含设备码的归档。修法是把静态库加进那个谓词(归档规则本来就消费 `lu.objects`);判据 e2e 608 断言 `ar t` 的成员表 —— 空档案也会成功 | +| 2026-09-05 | ⚠️⚠️ T5.1 的 CUDA lane 在本机**四维矩阵无解**,而链路本身全通 | 轴 → 收窄的 glob → 设备源清单 → 规则包 → action → **进静态库归档** → 链接,整条链实测走通,48 个设备目标已产出。挡住的是一个与 mcpp 无关的四维矩阵:①CCCL 2.x(12.9 线)的 `cub::LoadDirectWarpStriped` 少一个重载;②CCCL 3.3(13.3 线)同样不匹配;③补进 CCCL 3.2(13.2 线)后换成 **clang 编不动 libcu++**(`string_view` 的推导指引只允许 `__host__ __device__`、`block_load.cuh` 要 placement new);④走 nvcc 则 12.9 撞 glibc 2.44、13.3 撞驱动 12.4。⇒ 结论是**载荷矩阵**,不是设计 | +| 2026-09-05 | ⚠️ clang 路线要带 NVIDIA 自己的 libc++ 逃生开关 | 设备单元只要 include ``,`crt/host_defines.h:67` 就以 `"libc++ is not supported on x86 system"` 停下 —— 守卫是 `__CUDACC__ && _LIBCPP_VERSION`,而 clang 编 CUDA 时自己就定义 `__CUDACC__`,于是一条**写给 nvcc 宿主 pass** 的拒绝落到了这条路线上。`-D_ALLOW_UNSUPPORTED_LIBCPP` 只在 clang 路线传。⚠️ 示例自己的 kernel 一直没暴露它:**裸 kernel 一个工具包头都不 include**,示例测的是接线不是头文件 | +| 2026-09-05 | ⚠️ 未提交 `xim:cuda-cccl@13.2.75` | 配方改动做好并解析通过,但它是为一个没走通的用例加的,而索引的 CUDA 线策略是「12.9 + 13.3 两条」;单加一个 13.2 的 cccl 不自洽。撤回 | +| 2026-09-05 | ⚠️ **layer 不能选择依赖,feature 可以** | `[target.'cfg(accelerator = "cuda")'.dependencies]` 被引擎拒绝并说明理由:layer 由依赖图解析而来,用它选依赖会让依赖决定自己被问的问题。⇒ 一个库的设备后端拆两半:**依赖挂 feature,源文件挂 accel 轴** | +| 2026-09-05 | ⚠️ 设备编译要显式指名 CCCL 载荷,否则命中 `/usr/include/cub` | 与 T0.3 的 `cuda_runtime.h` 同一形状,第三次出现。且 12.x 的 `include/cub` 在 13.x 变成 `include/cccl/cub` | +| 2026-09-05 | e2e 317 的等待窗从 2s 放宽到 5s | 到达「五次短失败」下界最少要 1.25s(4×250ms 重启延迟 + 5×50ms 轮询),2s 窗只给每次 spawn 留 150ms;main 上 macOS **连续两次**在此失败而本分支同码两次通过 —— 判据由 runner 负载决定。5s 窗留 750ms | +| 2026-09-05 | e2e 602 声明 `requires: unix-shell`,并以 `MCPP_OFFLINE=1` 运行 | doctor 在 Windows 上整段不产出(载荷只有 linux 构建;Windows 工具包的上界是 `_MSC_VER` 区间,报告尚未读它);隔离 home 下 doctor 会把整套引导 + 工具链装进临时目录:实测 229s / 1.4 GB | +| 2026-09-05 | ⚠️ 核心改动:`--offline` 下跳过首次沙箱引导 | `load_or_init` 在空 home 里克隆索引、经 `xlings install` 装 ninja/patchelf,全部走网络,违反 `--offline`「绝不碰网络」的承诺。实测 offline 空 home 26s / 126 MB → 0.3s;e2e 604 带对照(已引导的 home 不提示);文档中英各补一句 | +| 2026-09-05 | C-1/C-2 不新增 `[[target]] kind = "device"` | 读了引擎:`mcpp::action` 已有四种角色(source/check/object/artifact),artifact 产物不进链接而 object 产物进链接,且 ninja 按路径连边 ⇒ 「不参与常规链接、由某条边消费的产物」就是 artifact 角色,「二次链接」就是以 artifact 为输入的 object 角色 action。再加一种 target kind 是同一个决定写第二遍 | +| 2026-09-05 | `accel` 进构建程序(`MCPP_ACCEL`)与 `cfg(accelerator)` | 后者的 `Ctx.accelerators` 字段**从未被写入**(声明了、文档了、没人填);e2e 605 第四段证明 `--no-accel` 下 layer 为空 | +| 2026-09-05 | ⚠️ `--accel/--no-accel` 进指纹并绕开 fast path | 实测:设备构建成功后 `mcpp build --no-accel` 报 `Finished in 0.00s` 并交回设备构建 —— 判据是 605 第四段先红后绿 | +| 2026-09-05 | ⚠️ 本机 e2e 168 红 | 已发布的 2026.9.5.1 同样红(musl gcc 13.3.0 载荷无 std module 源)⇒ 环境不是回归;CI 绿 | +| 2026-09-05 | ⚠️⚠️ 自构建的指纹目录中途漂移(`de4e07f` → `e45f24c7`) | 工作树 stash 后旧二进制也算出新值 ⇒ 输入在仓库外变了(后台 e2e 套件同时改写共享 `~/.mcpp`);我拿旧目录的二进制测了 606 近一小时。规则:**每次构建后用 `ls -t target/*/*/bin/mcpp` 重新取二进制,且开发期间不并行跑整套 e2e** | diff --git a/.agents/docs/2026-09-05-multi-device-verify.sh b/.agents/docs/2026-09-05-multi-device-verify.sh new file mode 100755 index 000000000..5f282c1c3 --- /dev/null +++ b/.agents/docs/2026-09-05-multi-device-verify.sh @@ -0,0 +1,209 @@ +#!/usr/bin/env bash +# Ecosystem verification for mcpp 2026.9.5.2, run inside a fresh xlings subos. +# +# ⚠️ EVERY ASSERTION CARRIES ITS OWN `|| fail`. A previous run of this kind was +# reported green because the transport dropped `set -euo pipefail` from the +# first line and the closing banner printed unconditionally. A script whose +# "pass" means "nothing failed" degrades, when that line is gone, into one whose +# "pass" means "it reached the last line", and the two read identically. +set -uo pipefail + +VER="${MCPP_VERIFY_VERSION:?set MCPP_VERIFY_VERSION}" +# The store path is the default and the point (see A); MCPP_VERIFY_BIN exists +# so this script can be rehearsed against a working-tree build before a release +# exists, and is never what a release is verified with. +STORE="${MCPP_VERIFY_BIN:-$HOME/.xlings/data/xpkgs/xim-x-mcpp/$VER/bin/mcpp}" + +fails=0 +fail() { printf 'ASSERT-FAIL: %s\n' "$1"; fails=$((fails + 1)); } +ok() { printf 'ok: %s\n' "$1"; } +have() { command -v "$1" >/dev/null 2>&1; } + +section() { printf '\n== %s ==\n' "$1"; } + +# ── A. the released binary is the one under test ──────────────────────────── +# +# ⭐ The store path, never the shim. `xlings install` has been observed to prune +# bare-name shims that mcpp itself installed, so a shim on PATH may resolve to +# an older version and the whole run would measure the wrong binary. +section "A. identity" +[ -x "$STORE" ] || fail "no released binary at $STORE" +got=$("$STORE" --version 2>&1 | head -1) +[ "$got" = "mcpp $VER" ] || fail "version is '$got', not 'mcpp $VER'" +[ "$got" = "mcpp $VER" ] && ok "$got from the store path" + +# ── B. the accelerator surface exists on all three verbs ──────────────────── +section "B. the device axis reaches run and test" +for verb in build run test; do + "$STORE" "$verb" --help 2>&1 | grep -q -- '--no-accel' \ + || fail "\`mcpp $verb --help\` does not mention --no-accel" +done +[ "$fails" -eq 0 ] && ok "build, run and test all take --accel/--no-accel" + +# ── C. a project that states a floor nothing satisfies is refused ─────────── +# +# The point is that the refusal happens BEFORE a compile, and names both values. +section "C. version floors are compared before compiling" +proj=$(mktemp -d) +mkdir -p "$proj/src" +cat > "$proj/mcpp.toml" <<'EOF' +[package] +name = "floor-probe" +version = "0.1.0" + +[language] +standard = "c++23" +modules = true +import_std = true + +[build] +sources = ["src/*.cpp"] + +[targets.floor-probe] +kind = "bin" +main = "src/main.cpp" +EOF +cat > "$proj/src/main.cpp" <<'EOF' +int main() { return 0; } +EOF +cat > "$proj/build.mcpp" <<'EOF' +import std; +import mcpp; +int main() { + mcpp::fact("probe.quantity", "1.0"); + mcpp::floor("probe.quantity >= 9.0"); + return 0; +} +EOF +out=$(cd "$proj" && "$STORE" build 2>&1) +printf '%s\n' "$out" | grep -q "requires probe.quantity >= 9.0" \ + || fail "an unmet floor did not refuse the build (output: $(printf '%s' "$out" | tail -3 | tr '\n' ' '))" +printf '%s\n' "$out" | grep -q "and this machine has 1.0" \ + || fail "the refusal did not state the measured value" +printf '%s\n' "$out" | grep -qi "compiling floor-probe" \ + && fail "the build compiled before the floor was compared" +printf '%s\n' "$out" | grep -q "requires probe.quantity >= 9.0" && ok "an unmet floor refuses before compiling, naming both values" + +# ── C'. the control: a floor that IS met does not refuse ──────────────────── +sed -i 's/probe.quantity >= 9.0/probe.quantity >= 0.5/' "$proj/build.mcpp" +out=$(cd "$proj" && "$STORE" build 2>&1) +printf '%s\n' "$out" | grep -q "Finished" \ + || fail "a satisfied floor still refused the build (output: $(printf '%s' "$out" | tail -3 | tr '\n' ' '))" +printf '%s\n' "$out" | grep -q "Finished" && ok "a satisfied floor builds — the check measures the floor, not the machine" + +# ── D. a constrained glob narrows, and --no-accel excludes it ─────────────── +# +# ⚠️ THE BUILD PROGRAM WRITES A FILE RATHER THAN PRINTING. mcpp shows a build +# program's stdout only when it exits non-zero, so a probe that printed its +# answer would be invisible on every successful build — and a grep that never +# matches is indistinguishable from one that matches nothing. +section "D. constrained source globs" +proj2=$(mktemp -d) +mkdir -p "$proj2/src/kernels" +cat > "$proj2/mcpp.toml" <<'EOF' +[package] +name = "glob-probe" +version = "0.1.0" +accelerators = ["cuda"] + +[language] +standard = "c++23" +modules = true +import_std = true + +[build] +accel = "cuda12.9+{sm_89}" +sources = [ + "src/*.cpp", + { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }, +] + +[targets.glob-probe] +kind = "bin" +main = "src/main.cpp" +EOF +cat > "$proj2/src/main.cpp" <<'EOF' +extern "C" int answer(); +int main() { return answer() == 7 ? 0 : 1; } +EOF +cat > "$proj2/src/answer.cpp" <<'EOF' +extern "C" int answer() { return 7; } +EOF +cat > "$proj2/src/kernels/k.cu" <<'EOF' +// Never compiled here: no rule package is imported, so a build that hands this +// file to the engine's compile rules is a build that narrowed wrongly. +#error "the device glob must not reach the engine's compile rules" +EOF +cat > "$proj2/build.mcpp" <<'EOF' +import std; +import mcpp; +int main() { + mcpp::rerun_if_env_changed("MCPP_PROBE_OUT"); + const char* out = std::getenv("MCPP_PROBE_OUT"); + if (out && *out) { + std::ofstream f(out, std::ios::trunc); + f << "DEVICE_SOURCES=[" << mcpp::device_sources() << "]\n"; + f << "ACCEL=[" << mcpp::accel() << "]\n"; + } + return 0; +} +EOF +probe_out="$proj2/probe.txt" +out=$(cd "$proj2" && MCPP_PROBE_OUT="$probe_out" "$STORE" build 2>&1) +printf '%s\n' "$out" | grep -q "Finished" \ + || fail "the device build failed (output: $(printf '%s' "$out" | tail -4 | tr '\n' ' '))" +grep -q "DEVICE_SOURCES=\[src/kernels/k.cu\]" "$probe_out" 2>/dev/null \ + || fail "the device source did not reach the build program (probe: $(cat "$probe_out" 2>/dev/null | tr '\n' ' '))" +grep -q "DEVICE_SOURCES=\[src/kernels/k.cu\]" "$probe_out" 2>/dev/null \ + && ok "a constrained glob reaches the build program as a device source" + +rm -f "$probe_out" +out=$(cd "$proj2" && MCPP_PROBE_OUT="$probe_out" "$STORE" build --no-accel 2>&1) +printf '%s\n' "$out" | grep -q "Finished" \ + || fail "--no-accel did not build (output: $(printf '%s' "$out" | tail -4 | tr '\n' ' '))" +grep -q "DEVICE_SOURCES=\[\]" "$probe_out" 2>/dev/null \ + || fail "--no-accel did not empty the device source list (probe: $(cat "$probe_out" 2>/dev/null | tr '\n' ' '))" +grep -q "ACCEL=\[\]" "$probe_out" 2>/dev/null \ + || fail "--no-accel did not empty the accel axis" +grep -q "DEVICE_SOURCES=\[\]" "$probe_out" 2>/dev/null \ + && ok "--no-accel excludes the constrained glob and still builds" + +out=$(cd "$proj2" && "$STORE" build --accel 'cuda12.9+{sm_80}' 2>&1) +printf '%s\n' "$out" | grep -q "does not cover" \ + || fail "an accel that does not cover the glob was not refused" +printf '%s\n' "$out" | grep -q "does not cover" && ok "an accel outside the constraint is refused naming both" + +# ── E. the engine carries no vendor tool or header name ──────────────────── +# +# ⚠️ ASSERTED ON THE BINARY, NOT ON `self doctor`'s SILENCE. An earlier +# revision asked whether `self doctor` mentions a toolkit, and that passes +# vacuously: the section it was watching for only ever printed when a CUDA +# toolkit was installed, and a fresh sandbox has none. Rehearsed against +# 2026.9.5.1 — which still had the reader — the check reported "ok". +# +# The strings are the discriminating ones: `crt/host_config.h` is the header +# the doctor used to read, and it appears twice in the 2026.9.5.1 binary and +# not at all once the reader moved into the rule package. The source-level +# property (no vendor name beside a probe launch anywhere in `src/`) is a unit +# test, `test_core_vendor_probes`, and belongs to CI rather than here. +section "E. the engine carries no vendor tool or header name" +vendor_hits=0 +for word in host_config cicc cudafe fatbinary nvidia-smi; do + # ⚠️ NO `|| echo 0`. `grep -c` PRINTS 0 and EXITS 1 when it matches nothing, + # so the fallback appended a second line and `[` saw "0\n0" — "integer + # expression expected", the assertion skipped, and the section still + # reported ok. Measured while writing this. + n=$(grep -a -c -- "$word" "$STORE" 2>/dev/null); rc=$? + [ "$rc" -gt 1 ] && { fail "could not read $STORE while looking for '$word'"; continue; } + [ "${n:-0}" -gt 0 ] && { fail "the engine binary contains '$word' ($n)"; vendor_hits=1; } +done +[ "$vendor_hits" -eq 0 ] && ok "no vendor tool or header name is in the engine binary" + +# ── verdict ──────────────────────────────────────────────────────────────── +printf '\n' +if [ "$fails" -eq 0 ]; then + printf 'ALL ASSERTIONS PASSED for mcpp %s\n' "$VER" + exit 0 +fi +printf '%d ASSERTION(S) FAILED for mcpp %s\n' "$fails" "$VER" +exit 1 diff --git a/CHANGELOG.md b/CHANGELOG.md index 28252a77d..ab936b25c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -5,22 +5,98 @@ ## [Unreleased] -### `mcpp self doctor` 报出设备编译器够不到自己的后端 - -一个 CUDA 工具包可以安装完整、就在 `PATH` 上,而设备代码仍然在第一个阶段失败。 -nvcc 以裸名调用 `cicc`、`cudafe++`、`ptxas`、`fatbinary`,依赖的是它自己从紧邻 -其二进制的 `nvcc.profile` 前置进来的一条 `PATH`。Debian 系打包把那个 profile -做成指向 `/etc` 的符号链接,于是任何替换了 `/etc` 的容器或沙箱都会移除它, -nvcc 沿用环境原有的 `PATH` 并报出 `sh: 1: cicc: not found` —— 这条消息既不提 -nvcc 也不提 profile,而工具包一样不缺,所以显而易见的检查全部通过。 - -判据取自 nvcc 自己:`--dryrun` 打印它将要运行的各个阶段与它将要使用的 `PATH`, -而不编译任何东西。mcpp 解析这份计划,逐个解析其中的裸名,报出第一个解析不到的。 -一次没有产生计划的 dryrun(没有 nvcc,或输出不是一份计划)不产生任何结论 —— -与宿主编译器上界那条检查同一条原则:一个够不到答案的探测不应当发明一个。 - -这是本轮把发布物放进沙箱验证时暴露出来的:示例在宿主上跑通,在 `--sandbox` -里失败,而唯一的差别就是那条 `PATH`。 +## [2026.9.5.2] — 2026-09-05 + +### ⭐⭐ 在编译任何东西之前比较机器的下界 + +有些机器事实**限定**了能为它构建什么,而忽略它们时,失败到得很晚。本轮的样本: +设备运行时不得新于它将运行其上的驱动;当它更新时,构建与链接都干净通过,程序在 +第一次分配处失败,消息里既不提工具包也不提驱动。 + +两个数字在编译任何东西之前都是可知的。mcpp 不去问厂商的工具要它们 —— +`tests/unit/test_runtime_contract` 禁止 `src/` 里出现厂商探针,而且这条规则是对的: +一个学会跑一家厂商探针的引擎会学会跑四家。所以数字以**声明**抵达: + +```toml +[[runtime.requirements]] +kind = "version-floor" +value = "cuda.driver >= 12.0" +``` + +`mcpp.build.version_floor` 只做比较,**这个文件里不出现任何厂商名字**: +`cuda.driver` 是流经的数据。第二种后端不需要改动它。 + +### ⭐ 探针通道:`mcpp::fact` / `mcpp::floor`(协议 v7) + +构建程序陈述它**测得**的事实与它**需要**的下界,引擎比较并在不满足时给出两侧取值 +(`version-floor-unmet`)。这是让 CUDA 探针得以整体离开 `src/` 的那条通道 —— +同样的读数现在由规则包产出,而它知道自己在跑哪个工具。 + +因此 **`mcpp self doctor` 的设备一节与 `mcpp.toolchain.devicehost` 一并删除**。 +它读得对(载荷优先于宿主、`crt/host_config.h` 的宿主编译器上界、`nvcc --dryrun` +的不可达阶段),但它不属于引擎。新增 `tests/unit/test_core_vendor_probes` +在剥掉注释的源码上陈述这条性质,并自带分母:枚举到的文件太少即判失败。 + +### 逐 glob 的加速器约束 + +`[build] sources` 的条目可以带上它面向的加速器: + +```toml +sources = [ + "src/*.cppm", + { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }, +] +``` + +构建按它收窄;`--no-accel` 整条排除;不覆盖它的 `--accel` 被拒并点名两侧 +(`accel-mismatch`);匹配为空的约束被拒并点名该 glob —— 空匹配是笔误或搬走了的 +目录,而不是空操作。 + +### `--accel` / `--no-accel` 现在也挂在 `run` 与 `test` 上 + +此前只有 `build` 有,实测后果是一个工程的 CPU-only 变体**能构建却不能运行**。 +两个动词接受同样的两个开关,给出任一个都绕开各自的快路径 —— 缓存的产物是按上一次 +构建的轴产的。 + +### 第二个编译器需要的两个答案 + +`MCPP_TOOLCHAIN_SYSROOT` 与 `MCPP_TOOLCHAIN_BINUTILS_DIR` 陈述 mcpp 传给**它自己** +那个编译器的 `--sysroot` 与 `-B`。规则包驱动一个 mcpp 并未解析的编译器时,那个 +编译器对环境一无所知:sub-OS 里 C 库不在 `/usr/include`,汇编器不在 `/usr/bin`, +于是它遇到的第一个 `#include` 就失败。`hipcc`、`-fsycl-host-compiler`、任何会编译 +自己产物的生成器都有同一个缺口。 + +`gcc::binutils_prefix_dir` 把 `-B` 的守卫收敛到一处,此前有三份副本,其中一份的 +注释写着它是另一份的镜像。 + +### Linux 上的动态构建程序 helper 改用 `DT_RPATH` + +RUNPATH 只对 helper **自己**的 needed 生效,于是一个在运行期打开宿主库的构建程序 +在下一跳失败:实测 `dlopen("/lib/libcuda.so.1")` 报 +`libdl.so.2: cannot open shared object file`,而持有它的目录就在 helper 的 RUNPATH 里。 +mcpp 链接的产物早就因为这个原因带 DT_RPATH。链接策略进 helper 的缓存身份, +旧 helper 会被重建而不是被重放。 + +### `--offline` 跳过首次使用的沙箱引导 + +`--offline` 承诺不碰网络,而 `load_or_init` 在空 home 里克隆索引、经 xlings 装 +ninja 与 patchelf。实测:空 home 下 26 秒 / 126 MB → 0.3 秒。 + +### `examples/09-cuda-kernel` 走两条路线,并拒绝它不能配的对 + +主路线是 clang(`-x cuda`):工程自己的编译器编设备单元,没有第二个宿主编译器、 +没有宿主编译器上界、没有 CUDA 的宿主头挡路。nvcc 是备用路线,并按名拒绝两对: +宿主编译器超出 `crt/host_config.h` 所述上界(实测 gcc 16 + nvcc 12.9 即便加了 +`-allow-unsupported-compiler` 也死在 gcc 自己的 `` 里),以及工具包 +旧于 C 库(12.9 的 `crt/math_functions.h` 为宿主重声明 C23 的 +`cospi`/`sinpi`/`rsqrt` 不带 `noexcept`,而 glibc 2.41+ 带)。 + +同一个缝下还有一份 CPU 实现,由 `cfg(not(accelerator = "cuda"))` 选中,于是 +`mcpp build --no-accel` 编它、`mcpp build` 编 `.cu`,两侧都不需要手写条件。 + +实测(RTX 4080,驱动 550.144.03 报 CUDA 12.4,LLVM 22.1.8):`mcpp run` 与 +`mcpp run --no-accel` 都打印 `12 24 36 48`,来自不同的产物目录,后者不含 +`cudaMalloc`。 ## [2026.9.5.1] — 2026-09-05 diff --git a/docs/00-getting-started.md b/docs/00-getting-started.md index ad81a716a..b1f5ea480 100644 --- a/docs/00-getting-started.md +++ b/docs/00-getting-started.md @@ -196,4 +196,4 @@ For the differences between the four modes and their artifact layouts, see [02 default/sole template, or run `mcpp new --list-templates ocornut.imgui@1.92.8`). - Explaining default decisions: `mcpp why [toolchain|runtime|deps]`; host capability checkup: `mcpp self doctor`; machine-readable resolution manifest: the build artifact `target///resolution.json`. -- Offline operation: `mcpp --offline` or `MCPP_OFFLINE=1` prevents index refreshes, downloads, and toolchain installation. +- Offline operation: `mcpp --offline` or `MCPP_OFFLINE=1` prevents index refreshes, downloads, and toolchain installation. In a home that has never been used it also skips the first-use sandbox bootstrap (index clone, ninja, patchelf), announces the skip once, and leaves the home un-bootstrapped; commands that need those tools report it. diff --git a/docs/05-mcpp-toml.md b/docs/05-mcpp-toml.md index 0876f2dac..579869619 100644 --- a/docs/05-mcpp-toml.md +++ b/docs/05-mcpp-toml.md @@ -215,6 +215,18 @@ the package/feature boundary, not on an individual target. > mcpp 2026.8.18.1 the two were byte-identical, so there was no spelling for > "nothing" and any file left under `src/` was swept in. +> **A `sources` entry may carry the accelerator it is for** (2026.9.5.2+): +> `{ glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }`. The glob +> joins the list like any other; the constraint decides whether it applies to a +> given build. It must match at least one file (an empty match is refused: it +> would leave nothing to compile for that device and say so only at the link). +> Under `--no-accel` the glob is left out, which is how one project yields its +> CPU-only variant. Under an `--accel` that does not cover the constraint the +> build is refused naming both (`accel-mismatch`). Device-kind files (`.cu`, +> `.hip`) the effective set matches are never compiled by the engine; they +> reach the build program as `MCPP_DEVICE_SOURCES`, where the rule package the +> project imports turns each into an `mcpp::action`. + ```toml [build] sources = ["src/**/*.cppm", "src/**/*.cpp"] # Source globs (default: src/**/*.{cppm,cpp,cc,c,S,s,asm}) @@ -1567,6 +1579,46 @@ error: capability 'gpu-blas' is provided by more than one package, and they The refusal reports `exclusive-capability` in `--format json` (chapter 11). +#### `version-floor` — needing more of the machine than it has + +Some facts about a machine bound what may be built for it, and the failure when +they are ignored arrives late: a program built against a runtime newer than the +driver it will meet links cleanly and fails at first use, naming neither side. + +A package states what it needs: + +```toml +[[runtime.requirements]] +kind = "version-floor" +value = "cuda.driver >= 12.0" +``` + +and a package that established a fact about this machine — at install time, +which is where probing belongs — states it: + +```toml +[runtime] +provides = ["cuda.driver=12.4"] +``` + +mcpp compares them when capabilities are bound and refuses before anything is +compiled, reporting `version-floor-unmet`: + +``` +error: `toolkitnew` requires cuda.driver >= 13.0, and this machine has 12.4. + stated by: driverfact +``` + +**No vendor vocabulary reaches the engine.** It reads a name, a relation and a +version; `cuda.driver` is data passing through, and a backend mcpp has never +heard of compares the same way. + +⚠️ **A floor nobody answered is silent.** A machine that never declared what it +has is not a machine that fails the floor — it is one nobody asked. Turning +"we do not know" into "no" is the failure mode this exists to avoid, and it is +asserted directly: `tests/e2e/603_version_floor.sh` builds a project whose floor +names something no package provides. + **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 diff --git a/docs/07-build-mcpp.md b/docs/07-build-mcpp.md index 34d6c91c9..d8f7c909b 100644 --- a/docs/07-build-mcpp.md +++ b/docs/07-build-mcpp.md @@ -55,6 +55,8 @@ is ignored, so diagnostics may be logged freely. | `mcpp:runner=` *(2026.8.19.2+)* | one argv token of the command that EXECUTES this build's artifact, when the host cannot. Emitted once per token, in order; the artifact path is appended (or substituted for `{}`). Reaches the **consumer**. ⚠️ Emit the executable as an ABSOLUTE path, and only **one** dependency may supply it | | `mcpp:link-script=` *(2026.8.19+)* | link with this **linker script** (`-T`; relative resolves against the package root, and the emitted path is absolute because the link runs in the build directory). Reaches the **consumer**, unlike `include-dir` — a board's memory layout is the one thing a consumer cannot write for itself | | `mcpp:warning=` *(2026.8.21.2+)* | say something to the user and **keep going**. The one directive that changes no compile line, no link line and no source set. Survives the build cache — see below | +| `mcpp:fact==` *(2026.9.5.2+)* | state something the program **established about the machine** (`cuda.driver=12.4`). Compared against floors before anything is compiled; see below | +| `mcpp:floor= >= ` *(2026.9.5.2+)* | state what this package **needs** of that quantity. Unmet ⇒ the build is refused with both values (`version-floor-unmet`); a floor nobody stated a fact for is silent | | `mcpp:rerun-if-changed=` | re-run `build.mcpp` when this file changes | | `mcpp:rerun-if-env-changed=` | re-run `build.mcpp` when this env var changes | @@ -150,6 +152,40 @@ there is nothing to do the build never reaches the `build.mcpp` stage — it als does not report which target it built or which sources it inferred. Touch a source and the advisory returns. +### The probe channel: `fact` / `floor` (2026.9.5.2+) + +A rule package is the thing that knows how to ask a machine what it has — +which library to open, which function to call — and the engine is the thing +that must not. So the package **measures** and the engine **compares**: + +```cpp +mcpp::fact("cuda.driver", "12.4"); // what this machine has +mcpp::floor("cuda.driver >= 12.0"); // what this package needs of it +``` + +Before anything is compiled, an unmet floor refuses the build and names the +quantity, both versions and who stated the fact; `mcpp why toolchain --format +json` classifies it as `reason: version-floor-unmet`. A floor for which nobody stated a fact is +**silent**: not knowing is not failing, and a refusal manufactured from +ignorance is the worse error. + +The failure this prevents is not visible at build time on its own. A program +built against a device runtime newer than the driver it will meet compiles +cleanly, links cleanly and fails at first use with a message naming neither +side. The rule package that resolved the runtime knows both numbers before +the first compile. + +⚠️ **Neither string means anything to the engine.** `cuda.driver` is data +flowing through; the engine reads a name, a relation and a version, and a +second backend needs no engine change. The spelling of a fact matches what a +package could also have declared statically in `[runtime] provides`, and a +floor matches `[[runtime.requirements]]` with `kind = "version-floor"`: the +two channels land in one list. + +⚠️ **A fact is cached with the program's other output** and replayed on a +cache hit. Declare what would change it — `rerun_if_changed` on the library +the version was read from — or the fact outlives the machine it described. + ### `runner` — how the artifact is executed (2026.8.19.2+) A board-support package knows the emulator, its machine model and its firmware @@ -199,6 +235,35 @@ package whose content is implementation-neutral. Asking follows whatever the compiler payload or through the runtime binding, and nothing has to look for it. +### Driving a second compiler: `toolchain_sysroot` / `toolchain_binutils_dir` (2026.9.5.2+) + +```cpp +const char* sr = mcpp::toolchain_sysroot(); // the `--sysroot` mcpp passes, or "" +const char* bu = mcpp::toolchain_binutils_dir(); // the dir mcpp names with `-B`, or "" +``` + +A rule package sometimes has to run a compiler mcpp did not resolve. `nvcc` +rejects a libc++ host compiler and fails inside GCC 16's ``, so a +CUDA rule package resolves a second host compiler from a declared payload; +`hipcc` and `-fsycl-host-compiler` pose the same question. + +That compiler starts knowing nothing about the environment it was placed in. +Under a sub-OS the C library is not at `/usr/include` and the assembler is not +at `/usr/bin`, so the first `#include` it reaches fails: + +``` +crt/host_config.h:218: fatal error: features.h: No such file or directory +``` + +These two answers are the flags mcpp passes to its own compiler for the same +target. Forwarding them — `--sysroot=` and `-B`, through whatever +the outer tool spells host options with — makes the second compiler see what +the first one sees. + +⚠️ **Not `sysroot_dir()`.** That answers a question about the target's *tier* +and is empty on a hosted target, which is exactly the case this pair exists +for. Either of these two is empty when mcpp passes no such flag. + ### Finding an `[xlings.workspace]` payload: `xpkg_dir` (2026.8.19+) `dep_dir` answers for **mcpp** dependencies. An xlings package is a different @@ -496,6 +561,10 @@ The running program receives the build context as `MCPP_*` variables | `MCPP_TARGET_ENV` *(0.0.100+)* | `mcpp::target_env()` | the target's env segment (`gnu`/`musl`/`msvc`); empty string when the triple has none (macOS) | | `MCPP_HOST` | `mcpp::host()` | the host triple | | `MCPP_PROFILE` | `mcpp::profile()` | effective profile name (`dev`/`release`/…) | +| `MCPP_TOOLCHAIN_SYSROOT` *(2026.9.5.2+)* | `mcpp::toolchain_sysroot()` | the `--sysroot` mcpp passes to its own compiler; empty when it passes none. For a rule package that runs a **second** compiler — see "Driving a second compiler" above | +| `MCPP_TOOLCHAIN_BINUTILS_DIR` *(2026.9.5.2+)* | `mcpp::toolchain_binutils_dir()` | the directory mcpp names with `-B`; empty when it names none (a musl or MinGW payload brings its own assembler and linker) | +| `MCPP_ACCEL` *(2026.9.5.2+)* | `mcpp::accel()` | the device axis of this build, resolved — `--accel` / `--no-accel` over `[build] accel` — in the wire form `cuda12.9+{sm_89} ptx>=89`; empty when the build asks for no accelerator. A rule package derives its own flags (`-gencode`, `--offload-arch`) from it, so the architecture set is written once, in the manifest. The same value feeds the `cfg(accelerator = "…")` layer key | +| `MCPP_DEVICE_SOURCES` *(2026.9.5.2+)* | `mcpp::device_sources()` | the device-kind sources (`.cu`, `.hip`, …) the package's effective `sources` match, package-root-relative, one per line; empty when there are none. The engine compiles none of them — the rule package this program imports turns each into an `mcpp::action`. Already narrowed: a `{ glob, accel }` entry the build does not cover contributes nothing, so `--no-accel` yields an empty list | | `MCPP_OUT_DIR` | `mcpp::out_dir()` | a writable scratch/output dir owned by mcpp | | `MCPP_MANIFEST_DIR` | `mcpp::manifest_dir()` | the package root (= CWD) | | `MCPP_FEATURE_` | `mcpp::has_feature("name")` | set to `1` per active feature (same `` sanitization as the `MCPP_FEATURE_` compile macro) | diff --git a/docs/11-machine-output.md b/docs/11-machine-output.md index 39f8eae0d..1b0cd536e 100644 --- a/docs/11-machine-output.md +++ b/docs/11-machine-output.md @@ -378,6 +378,8 @@ a program classifying the outcome reads `reason`: | `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 | +| `version-floor-unmet` | a package requires more of the machine than the machine was declared to have | +| `accel-mismatch` | a `[build] sources` entry is constrained to a device set this build does not cover | | `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/20-accelerators.md b/docs/20-accelerators.md index bab3b090f..82c1ba78c 100644 --- a/docs/20-accelerators.md +++ b/docs/20-accelerators.md @@ -67,55 +67,47 @@ The division is deliberate. mcpp owns the graph, the artifact's identity and the set of architectures; a vendor's flag spelling, its architecture syntax and its host-compiler requirements belong to the rule. -## The host compiler a device compiler will accept - -nvcc refuses host compilers newer than a bound it states in its own -`crt/host_config.h`, and mcpp's toolchain payload is frequently newer than that -bound. Because mcpp supplies the host compiler, it can report the pairing -before anything is compiled: - -``` -$ mcpp self doctor - Checking device toolkit -warning: cuda will refuse this host compiler: gcc 13 exceeds the bound of 12 - stated in /usr/include/crt/host_config.h. -``` - -The bound is read from the toolkit rather than tabulated in mcpp, so a toolkit -mcpp has never seen still answers, and a header mcpp cannot parse yields no -bound and therefore no claim. - -This is reported rather than enforced: a project that compiles no device code -is unaffected by an incompatible pair. - -## Whether the device compiler can reach its own back-end - -A toolkit can be installed, complete and on `PATH` and still fail at its first -stage. nvcc runs `cicc`, `cudafe++`, `ptxas` and `fatbinary` as bare names, on -a `PATH` it prepends itself from an `nvcc.profile` beside its own binary. On -Debian-family packaging that profile is a symlink into `/etc`, so a container -or sandbox that replaces `/etc` removes it. nvcc then keeps the ambient `PATH` -and reports: - -``` -sh: 1: cicc: not found -``` - -The message names neither nvcc nor the profile, and nothing about the toolkit -is missing, so the obvious checks all pass. `mcpp self doctor` asks nvcc for -its plan instead of assuming one: - -``` -$ mcpp self doctor - Checking device toolkit -warning: nvcc cannot reach its own back-end: it invokes 'cicc' by name, and - that name does not resolve on the search path it states. -``` - -The plan comes from `nvcc --dryrun`, which prints the stages and the `PATH` -nvcc will use without compiling anything. A dryrun that produces no plan -- -there is no nvcc, or the output is not one -- yields no finding, because a -probe that reaches no answer must not invent one. +## What the rule package reports before the first compile + +Three things go wrong late with a device toolkit, and none of them is a fact +about the build graph. They are read and reported by the **rule package** that +drives the tools -- `examples/09-cuda-kernel/rules-cuda` shows each one -- and +the engine owns none of them (`tests/unit/test_core_vendor_probes.cpp` holds +that line, so a second backend never grows a second copy inside mcpp). + +**The host compiler a device compiler will accept.** nvcc refuses host +compilers newer than a bound it states in its own `crt/host_config.h`, and +mcpp's toolchain payload is frequently newer than that bound. On the nvcc +route the rule reads the bound from the toolkit it resolved -- a payload +before the host, because a toolkit installed through xlings is the one the +build uses and is usually the newer one (a 12.9 payload states `gcc <= 14` +where a distribution's CUDA 12.0 states `gcc <= 12`) -- and says which +compiler it chose and why, through `mcpp::warning`. The primary route has no +such bound: `clang -x cuda` is its own host compiler. + +**Whether the device compiler can reach its own back-end.** A toolkit can be +installed, complete and on `PATH` and still fail at its first stage: nvcc +runs `cicc`, `cudafe++`, `ptxas` and `fatbinary` as bare names on a `PATH` it +prepends from an `nvcc.profile` beside its own binary, and a sandbox that +replaces `/etc` removes a Debian-packaged profile. The rule asks nvcc for its +plan (`nvcc --dryrun`) rather than assuming one, resolves each stage, and +names the first one that does not resolve together with the payload that +provides it. A dryrun that produces no plan yields no finding. + +**Whether the driver is new enough for the runtime.** A device runtime must +not be newer than the driver it runs against; when it is, the build compiles +and links cleanly and fails at the first allocation with *"CUDA driver +version is insufficient for CUDA runtime version"*. The rule reads the +driver's version through the driver's own library (reached through the +sentinel package, never through `/usr/lib`) and states it as a fact; it states +the floor its runtime needs; and the engine compares the two before anything +is compiled -- see the probe channel in [07 — build.mcpp](07-build-mcpp.md). +The engine reads a name, a relation and a version; `cuda.driver` is data +flowing through. + +Reported rather than enforced where a wrong answer would cost more than none: +a machine with no rule package in its project has nothing vendor-specific to +say and says nothing, and a probe that reaches no answer invents none. ## Declaring what a build targets @@ -239,10 +231,12 @@ combinators. ## Two boundaries worth stating -**`--accel` is a `build` option**, alongside `--static` and `--toolchain`, and -is not repeated on `run`, `test` or `pack`. Those read `[build] accel` from the -manifest like every other build input; the flag exists for overriding one -build, which is the case `build` covers. +**`--accel` and `--no-accel` are `build`, `run` and `test` options** (run and +test from 2026.9.5.2), as `--target` and `--profile` are. `pack` reads +`[build] accel` from the manifest like every other build input. The flag was a +`build`-only option at first, and the measured consequence was a project whose +CPU-only variant could be built and not run: `mcpp build --no-accel` produced +it, and `mcpp run` handed back the device build. **`mcpp pack` does not emit the `accel` field.** It could write whatever the manifest declared, and that is exactly why it does not: the field states what an diff --git a/docs/zh/00-getting-started.md b/docs/zh/00-getting-started.md index e28bf6163..745f33f16 100644 --- a/docs/zh/00-getting-started.md +++ b/docs/zh/00-getting-started.md @@ -183,4 +183,4 @@ mcpp pack --mode self-contained # 打包 loader、libc 与依赖 `mcpp new --list-templates ocornut.imgui@1.92.8`)。 - 解释默认决策:`mcpp why [toolchain|runtime|deps]`;主机能力体检:`mcpp self doctor`; 机器可读解析清单:构建产物 `target///resolution.json`。 -- 离线运行:`mcpp --offline` 或 `MCPP_OFFLINE=1` 可阻止索引刷新、下载和工具链安装。 +- 离线运行:`mcpp --offline` 或 `MCPP_OFFLINE=1` 可阻止索引刷新、下载和工具链安装。在从未使用过的 home 中,它同时跳过首次使用时的沙箱引导(索引克隆、ninja、patchelf),只提示一次,并让该 home 保持未引导状态;需要这些工具的命令会自行报告。 diff --git a/docs/zh/05-mcpp-toml.md b/docs/zh/05-mcpp-toml.md index c57452843..e5a29653d 100644 --- a/docs/zh/05-mcpp-toml.md +++ b/docs/zh/05-mcpp-toml.md @@ -201,6 +201,15 @@ mcpp 刻意不在一次构建里把同一个共享源编译成两份:一个源 > `src/` 下剩下的任何文件都会被扫进来。 +> **`sources` 的条目可以带上它所面向的加速器**(2026.9.5.2+): +> `{ glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }`。glob 与其它条目一样 +> 进入列表;约束决定它是否适用于某一次构建。它必须至少匹配一个文件(空匹配会被拒绝: +> 那会让这个设备无东西可编,而只在链接时才说话)。`--no-accel` 下该 glob 被排除, +> 一个工程由此产出它的 CPU-only 变体。`--accel` 未覆盖该约束时构建被拒并给出两侧 +> (`accel-mismatch`)。有效集合匹配到的设备类源文件(`.cu`、`.hip`)引擎从不编译; +> 它们以 `MCPP_DEVICE_SOURCES` 到达构建程序,由工程引入的规则包把每一个变成一条 +> `mcpp::action`。 + ```toml [build] sources = ["src/**/*.cppm", "src/**/*.cpp"] # 源文件 glob(默认: src/**/*.{cppm,cpp,cc,c,S,s,asm}) @@ -1347,6 +1356,42 @@ error: capability 'gpu-blas' is provided by more than one package, and they 该拒绝在 `--format json` 里报 `exclusive-capability`(见第 11 章)。 +#### `version-floor` —— 对机器的要求高于它所有 + +有些关于机器的事实约束着能为它构建什么,而忽略它们时的失败来得很晚: +一个针对比它将遇到的驱动更新的运行时构建出来的程序,**干净地链接**, +在第一次使用时失败,而消息里两侧都没有。 + +包声明它需要什么: + +```toml +[[runtime.requirements]] +kind = "version-floor" +value = "cuda.driver >= 12.0" +``` + +而某个在**安装期**(探测该发生的地方)确立了机器某项事实的包,声明它: + +```toml +[runtime] +provides = ["cuda.driver=12.4"] +``` + +mcpp 在绑定 capability 时比较二者,并在任何东西被编译之前拒绝, +报 `version-floor-unmet`: + +``` +error: `toolkitnew` requires cuda.driver >= 13.0, and this machine has 12.4. + stated by: driverfact +``` + +**没有任何厂商词汇抵达引擎。** 它读到的是一个名字、一个关系和一个版本; +`cuda.driver` 是流过的数据,一个 mcpp 从未听说过的后端比较方式完全相同。 + +⚠️ **没人回答的下界是沉默的。** 一台从未声明自己有什么的机器,不是「未满足下界」的机器, +而是「没人问过」的机器。把「我们不知道」变成「不行」正是这个机制要避免的失败, +并且有直接判据:`tests/e2e/603_version_floor.sh` 会构建一个下界指向无人提供之物的工程。 + **它是关于这个包自己的符号的声明**,所以一条指向本包并不提供的能力的条目会被报为 schema 警告:那里没有可独占的东西。而无人声明独占的能力行为完全不变 —— 两个 BLAS 实现照常共存,既有的「两个或更多、未 pin」报错也仍然只在**有人 require** 该能力时出现。 diff --git a/docs/zh/07-build-mcpp.md b/docs/zh/07-build-mcpp.md index 17dc8482e..50e906c18 100644 --- a/docs/zh/07-build-mcpp.md +++ b/docs/zh/07-build-mcpp.md @@ -52,6 +52,8 @@ mcpp build # 编译 + 运行 build.mcpp,然后构建工程 | `mcpp:runner=` *(2026.8.19.2+)* | 执行本次构建产物的命令的**一个 argv token**(宿主跑不了它时)。一个 token 一次调用、按顺序;产物路径会被追加(或替换 `{}`)。**到达消费者**。⚠️ 可执行文件要发**绝对路径**,且**只能有一个**依赖提供它 | | `mcpp:link-script=` *(2026.8.19+)* | 用这个**链接脚本**链接(`-T`;相对路径按包根解析,发出的是绝对路径,因为链接是在构建目录里跑的)。与 `include-dir` 不同,它**到达消费者** —— 板子的内存布局恰恰是消费者写不出来的那一项 | | `mcpp:warning=` *(2026.8.21.2+)* | 对用户说一句话并**继续**。唯一一条不改变编译行、链接行与源码集的指令。它**穿过构建缓存** —— 见下 | +| `mcpp:fact==` *(2026.9.5.2+)* | 陈述程序**测得的机器事实**(`cuda.driver=12.4`)。在编译任何东西之前与 floor 比较;见下 | +| `mcpp:floor= >= ` *(2026.9.5.2+)* | 陈述本包对该量的**下界**。不满足 ⇒ 构建被拒并给出两侧取值(`version-floor-unmet`);没有人陈述事实的下界保持沉默 | | `mcpp:rerun-if-changed=` | 该文件变化时重跑 `build.mcpp` | | `mcpp:rerun-if-env-changed=` | 该环境变量变化时重跑 `build.mcpp` | @@ -137,6 +139,32 @@ mcpp 在每次命中时重放它。 `build.mcpp` 阶段 —— 它同样不会报告构建了哪个目标、推断了哪些源码。touch 一下源码,提示 就回来了。 +### 探针通道:`fact` / `floor`(2026.9.5.2+) + +规则包是知道「怎么问机器它有什么」的那一方 —— 打开哪个库、调用哪个函数; +引擎则是不该知道的那一方。于是由包来**测量**,由引擎来**比较**: + +```cpp +mcpp::fact("cuda.driver", "12.4"); // 这台机器有什么 +mcpp::floor("cuda.driver >= 12.0"); // 本包需要它至少多新 +``` + +在编译任何东西之前,未满足的下界会拒绝构建,并点名该量、两侧版本与陈述事实的 +包;`mcpp why toolchain --format json` 把它归类为 `reason: version-floor-unmet`。没有人陈述事实的下界 +**保持沉默**:不知道不等于不满足,由无知制造的拒绝是更坏的错误。 + +它防住的失败在构建期本身看不见:针对比驱动更新的设备运行时构建的程序,编译 +干净、链接干净,到第一次使用才失败,而消息两边都不点名。解析了该运行时的规则包 +在第一次编译之前就知道两个数字。 + +⚠️ **两个字符串对引擎都没有意义。** `cuda.driver` 是流过引擎的数据;引擎读到的 +是一个名字、一个关系、一个版本,第二个后端不需要改引擎。事实的拼法与包在 +`[runtime] provides` 里静态声明的一致,下界与 `[[runtime.requirements]]` 的 +`kind = "version-floor"` 一致:两条通道落进同一张表。 + +⚠️ **事实随程序的其它输出一起进缓存**,命中时被回放。要声明什么会改变它 —— +对读出版本的那个库 `rerun_if_changed` —— 否则事实会比它描述的机器活得更久。 + ### `runner` —— 产物的执行方式(2026.8.19.2+) 板级支持包知道模拟器、机器型号和固件模式,也知道模拟器**在哪** —— 而静态 manifest @@ -177,6 +205,31 @@ const char* sr = mcpp::sysroot_dir(); // 目标的 C 库根目录,没有则 宿主目标上 `sysroot_dir()` 为空:那里 C 库随编译器载荷或运行时绑定而来,没人需要找它。 +### 驱动第二个编译器:`toolchain_sysroot` / `toolchain_binutils_dir`(2026.9.5.2+) + +```cpp +const char* sr = mcpp::toolchain_sysroot(); // mcpp 传的 `--sysroot`,没有则为 "" +const char* bu = mcpp::toolchain_binutils_dir(); // mcpp 用 `-B` 指的目录,没有则为 "" +``` + +规则包有时必须运行一个 **mcpp 并未解析**的编译器。`nvcc` 拒绝 libc++ 宿主编译器, +在 GCC 16 的 `` 上失败,因此 CUDA 规则包要从声明的载荷里另选一个宿主 +编译器;`hipcc` 与 `-fsycl-host-compiler` 面对同一个问题。 + +那个编译器对自己被放进的环境一无所知。在 sub-OS 里 C 库不在 `/usr/include`,汇编器 +不在 `/usr/bin`,于是它遇到的第一个 `#include` 就失败: + +``` +crt/host_config.h:218: fatal error: features.h: No such file or directory +``` + +这两个答案就是 mcpp 为同一目标传给自己那个编译器的开关。把它们转发过去 +——`--sysroot=<值>` 与 `-B<值>`,经外层工具的宿主选项拼法——第二个编译器就看到 +第一个看到的东西。 + +⚠️ **不是 `sysroot_dir()`。** 那个回答的是目标**档位**的问题,在宿主目标上为空, +而宿主目标恰恰是这一对存在的场合。mcpp 不传某个开关时,对应的那个为空串。 + ### 找到 `[xlings.workspace]` 的载荷:`xpkg_dir`(2026.8.19+) `dep_dir` 回答的是 **mcpp** 依赖。xlings 包是另一个命名空间、另一套 store 布局, @@ -434,6 +487,10 @@ mcpp 会把它自己构建时用的**同一份** std 模块暂存过来,缓存 | `MCPP_TARGET_ENV` *(0.0.100+)* | `mcpp::target_env()` | 目标的 env 段(`gnu`/`musl`/`msvc`);三元组无 env 段(macOS)时为空串 | | `MCPP_HOST` | `mcpp::host()` | 宿主三元组 | | `MCPP_PROFILE` | `mcpp::profile()` | 生效 profile 名(`dev`/`release`/…) | +| `MCPP_TOOLCHAIN_SYSROOT` *(2026.9.5.2+)* | `mcpp::toolchain_sysroot()` | mcpp 传给自己那个编译器的 `--sysroot`;不传时为空串。供运行**第二个**编译器的规则包使用 —— 见上文「驱动第二个编译器」 | +| `MCPP_TOOLCHAIN_BINUTILS_DIR` *(2026.9.5.2+)* | `mcpp::toolchain_binutils_dir()` | mcpp 用 `-B` 指的目录;不指时为空串(musl 与 MinGW 载荷自带汇编器与链接器) | +| `MCPP_ACCEL` *(2026.9.5.2+)* | `mcpp::accel()` | 本次构建的设备轴,已解析 —— `--accel` / `--no-accel` 优先于 `[build] accel` —— 线上形态 `cuda12.9+{sm_89} ptx>=89`;不要加速器时为空串。规则包从它推导自己的开关(`-gencode`、`--offload-arch`),架构集合因此只在 manifest 写一次。同一个值也喂给 `cfg(accelerator = "…")` 这个 layer 键 | +| `MCPP_DEVICE_SOURCES` *(2026.9.5.2+)* | `mcpp::device_sources()` | 本包有效 `sources` 匹配到的设备类源文件(`.cu`、`.hip`…),相对包根,一行一个;没有时为空串。引擎一个都不编译 —— 由本程序引入的规则包把每一个变成一条 `mcpp::action`。已经过收窄:构建未覆盖的 `{ glob, accel }` 条目贡献为空,因此 `--no-accel` 得到空列表 | | `MCPP_OUT_DIR` | `mcpp::out_dir()` | mcpp 提供的可写输出/暂存目录 | | `MCPP_MANIFEST_DIR` | `mcpp::manifest_dir()` | 包根(= CWD) | | `MCPP_FEATURE_` | `mcpp::has_feature("name")` | 每个活跃 feature 置 `1`(`` 消毒规则与 `MCPP_FEATURE_` 编译宏一致) | diff --git a/docs/zh/11-machine-output.md b/docs/zh/11-machine-output.md index d39f0b496..06093d080 100644 --- a/docs/zh/11-machine-output.md +++ b/docs/zh/11-machine-output.md @@ -332,6 +332,8 @@ mcpp why toolchain [--target ] [--toolchain ] --format json | `layer-requirement` | 某个包要求的层,解析没有给出 | | `layer-ordering` | 五层叠不起来 | | `exclusive-capability` | 一个能力有多个提供者,而其中至少一个声明了独占 | +| `version-floor-unmet` | 一个包对机器的要求高于机器被声明拥有的 | +| `accel-mismatch` | 一条 `[build] sources` 条目被约束到本次构建未覆盖的设备集合 | | `other` | 一处还没有被命名的拒绝分支 | ⚠️ **只要问题被回答了就退 0,包括答案是「拒绝」。** 「它能不能构建,不能的话 diff --git a/docs/zh/20-accelerators.md b/docs/zh/20-accelerators.md index a056129c6..e3415dd80 100644 --- a/docs/zh/20-accelerators.md +++ b/docs/zh/20-accelerators.md @@ -51,50 +51,35 @@ tarball 已经发出去了。设备源文件必须被显式点名。 这个划分是刻意的。mcpp 拥有构建图、产物身份与架构集合;厂商的 flag 拼法、 架构语法与宿主编译器要求属于规则包。 -## 设备编译器接受哪些宿主编译器 - -nvcc 拒绝比它在自己的 `crt/host_config.h` 中声明的上界更新的宿主编译器, -而 mcpp 的工具链载荷常常比那个上界更新。由于宿主编译器由 mcpp 提供, -它可以在任何编译发生之前报告这个配对: - -``` -$ mcpp self doctor - Checking device toolkit -warning: cuda will refuse this host compiler: gcc 13 exceeds the bound of 12 - stated in /usr/include/crt/host_config.h. -``` - -上界是从工具包读出的,不是抄在 mcpp 里的表,因此一个 mcpp 从未见过的工具包同样能作答; -而一个 mcpp 无法解析的头文件产生不出上界,也就不产生任何断言。 - -这是报告而非强制:一个不编译任何设备代码的工程,不受不兼容配对的影响。 - -## 设备编译器能否够到自己的后端 - -一个工具包可以安装完整、就在 `PATH` 上,却仍然在第一个阶段失败。 -nvcc 以裸名调用 `cicc`、`cudafe++`、`ptxas` 与 `fatbinary`,依赖的是它自己 -从紧邻其二进制的 `nvcc.profile` 前置进来的一条 `PATH`。在 Debian 系的打包里, -那个 profile 是指向 `/etc` 的符号链接,于是任何替换了 `/etc` 的容器或沙箱都会移除它。 -nvcc 随即沿用环境里原有的 `PATH`,并报出: - -``` -sh: 1: cicc: not found -``` - -这条消息既没有提到 nvcc,也没有提到 profile,而工具包本身一样不缺, -于是所有显而易见的检查都会通过。`mcpp self doctor` 因此去问 nvcc 要它的计划, -而不是假设一份: - -``` -$ mcpp self doctor - Checking device toolkit -warning: nvcc cannot reach its own back-end: it invokes 'cicc' by name, and - that name does not resolve on the search path it states. -``` - -计划来自 `nvcc --dryrun` —— 它打印各个阶段与 nvcc 将要使用的 `PATH`, -而不编译任何东西。一次没有产生计划的 dryrun(没有 nvcc,或输出不是一份计划) -不产生任何结论:一个够不到答案的探测不应当发明一个。 +## 规则包在第一次编译之前报告的事 + +设备工具包有三件事出错得很晚,而没有一件是关于构建图的事实。它们由驱动这些工具的 +**规则包**读取并报告 —— `examples/09-cuda-kernel/rules-cuda` 逐一演示 —— 引擎一件都 +不拥有(`tests/unit/test_core_vendor_probes.cpp` 守住这条线,于是第二个后端不会在 +mcpp 里长出第二份拷贝)。 + +**设备编译器接受哪些宿主编译器。** nvcc 拒绝比它在自己的 `crt/host_config.h` 中声明 +的上界更新的宿主编译器,而 mcpp 的工具链载荷往往比那个上界更新。在 nvcc 路线上,规则 +从它解析到的工具包读出上界 —— 载荷先于宿主,因为经 xlings 安装的工具包才是构建会用 +的那个,通常也是更新的那个(12.9 载荷写着 `gcc <= 14`,发行版的 CUDA 12.0 写着 +`gcc <= 12`)—— 并通过 `mcpp::warning` 说出它选了哪个编译器、为什么。主路线没有这条 +上界:`clang -x cuda` 自己就是宿主编译器。 + +**设备编译器能否够到自己的后端。** 工具包可以装好、完整、在 `PATH` 上,却在第一阶段 +就失败:nvcc 以裸名调用 `cicc`、`cudafe++`、`ptxas`、`fatbinary`,靠的是它从自己二进制 +旁边的 `nvcc.profile` 前置进来的一条 `PATH`,而替换了 `/etc` 的沙箱会拿走 Debian 打包的 +那个 profile。规则向 nvcc 要它的计划(`nvcc --dryrun`)而不是假设一个,逐个解析各阶段, +点名第一个解析不到的以及提供它的载荷。产不出计划的 dryrun 不产生任何结论。 + +**驱动是否新到足以承载运行时。** 设备运行时不能比它将遇到的驱动更新;更新时,构建编译 +干净、链接干净,到第一次分配才以 *"CUDA driver version is insufficient for CUDA runtime +version"* 失败。规则经驱动自己的库(经由 sentinel 包够到,绝不经 `/usr/lib`)读出驱动 +版本并陈述为事实;陈述它的运行时需要的下界;引擎在编译任何东西之前比较两者 —— 见 +[07 — build.mcpp](07-build-mcpp.md) 的探针通道。引擎读到的是一个名字、一个关系、一个 +版本;`cuda.driver` 是流过引擎的数据。 + +凡是错答比不答更贵的地方都只报告不强制:工程里没有规则包的机器没有任何厂商相关的话 +要说,于是什么都不说;够不到答案的探针不发明答案。 ## 声明一次构建的目标 @@ -204,7 +189,9 @@ cxxflags = ["-DMYAPP_ROCM"] ## 两条值得写明的边界 -**`--accel` 是 `build` 的选项**,与 `--static`、`--toolchain` 同级, +**`--accel` 与 `--no-accel` 是 `build`、`run`、`test` 三者的选项**(run 与 test 自 2026.9.5.2 起),与 `--target`、`--profile` 同级;`pack` 与其它构建输入一样从 manifest 读 `[build] accel`。它起初只挂在 `build` 上,实测的后果是一个工程的 CPU-only 变体能构建却不能运行:`mcpp build --no-accel` 产出了它,而 `mcpp run` 交回的是设备构建。 + +**历史:`--accel` 曾只是 `build` 的选项**,与 `--static`、`--toolchain` 同级, 不在 `run`、`test`、`pack` 上重复。那些命令与读取任何其它构建输入一样, 从 manifest 读 `[build] accel`;这个 flag 的用途是覆盖单次构建,而那正是 `build` 覆盖的场景。 diff --git a/examples/09-cuda-kernel/README.md b/examples/09-cuda-kernel/README.md index 2853b43fd..8ef9ec57b 100644 --- a/examples/09-cuda-kernel/README.md +++ b/examples/09-cuda-kernel/README.md @@ -6,14 +6,17 @@ What this example demonstrates, and what it does not. ``` app/ - src/kernels/saxpy.cu the island: compiled by nvcc, never scanned, no BMI + src/kernels/saxpy.cu the island: a device translation unit, never scanned, + no BMI, and compiled only when the build asks for CUDA + src/cpu/saxpy.cpp the same interface implemented for the host, compiled + instead when it does not include/saxpy/saxpy.h the island's interface: extern "C", no std types src/app.cppm the seam: a module that turns the C interface back into a C++ one src/main.cpp an ordinary consumer, which imports the seam and never sees the header - build.mcpp names the sources and the architectures -rules-cuda/ a build-rule package that knows how to run nvcc + build.mcpp hands the device sources to the rule package +rules-cuda/ a build-rule package that knows how to compile them ``` Three properties are load-bearing. @@ -24,49 +27,178 @@ imports, never producing a BMI. Its header is classified as a header, so editing one still invalidates the fast path. **The island's interface is `extern "C"` and free of standard-library types.** -nvcc drives a host compiler that mcpp did not choose, so the two sides do not -share a C++ ABI and must not exchange anything that depends on one. The island -also uses no standard library itself, which keeps it from linking a second copy -of the C++ runtime into a program whose own copy came from mcpp's toolchain. +The device unit may be compiled by a compiler mcpp did not choose, so the two +sides do not share a C++ ABI and must not exchange anything that depends on +one. The island also uses no standard library itself, which keeps it from +linking a second copy of the C++ runtime into a program whose own copy came +from mcpp's toolchain. **The seam exists for backend substitution, not for the module boundary.** It -is the one place where the island underneath could become HIP or a CPU -fallback without any consumer of `app.saxpy` changing, and the one place a -`cfg(accelerator = ...)` section has to apply. Remove it and every importer -becomes backend-specific. +is the one place where the island underneath becomes a CPU implementation, or +could become HIP, without any consumer of `app.saxpy` changing, and the one +place a `cfg(accelerator = ...)` section has to apply. Remove it and every +importer becomes backend-specific. -## The rule package, and why nvcc's host compiler is its problem +## Two builds from one source tree -nvcc refuses host compilers newer than a bound it states in its own -`crt/host_config.h`, and mcpp's toolchain payload is routinely newer than that -bound. The rule reads the bound, selects a host compiler that satisfies it, and -says which one it chose: +The device axis is written once, in the manifest, and the source set follows +it: + +```toml +[build] +accel = "cuda12.9+{sm_89} ptx>=89" +sources = [ + "src/*.cppm", + "src/*.cpp", + { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }, +] + +[target.'cfg(not(accelerator = "cuda"))'.build] +sources = ["src/cpu/*.cpp"] +``` + +A glob may carry the accelerator it is for. `mcpp build` compiles the `.cu` +and the CPU file is absent; `mcpp build --no-accel` compiles the CPU file and +the `.cu` is absent — not excluded by a hand-written condition, but by the +constraint the glob states. The two land in different artifact directories +because the device axis is part of the build's identity, so switching between +them does not rebuild from scratch. + +An `--accel` that does not cover a constrained glob is refused before anything +is compiled, with `accel-mismatch` on the machine-readable channel. + +## Where the toolkit comes from + +The project names it: + +```toml +[xlings.workspace] +"xim:cuda-nvcc" = "12.9.86" +"xim:cuda-cudart" = "12.9.79" +``` + +These are payloads, so the version is the project's choice and not the +machine's. The rule package resolves them with `mcpp::xpkg_dir` and builds the +whole invocation from what it finds — the compiler, the include directories +and the library search paths. No path in this example is absolute, and a build +here touches nothing of the host's CUDA: ``` -example.rules.cuda: nvcc /usr/bin/nvcc with -ccbin /usr/bin/clang++-14 +$ mcpp build -v | grep -c '/usr/local/cuda\|/usr/bin/nvcc' +0 ``` -On the machine this example was verified on, the toolkit is CUDA 12.0 -(`__GNUC__ > 12` is refused, clang must be below 15) and mcpp's payload is gcc -16.1.0, so passing mcpp's own compiler through would fail. `mcpp self doctor` -reports the same pairing independently. +**The payload's headers have to be named.** nvcc adds +`/../include` by itself, and on the 12.x line that holds +`crt/` but not `cuda_runtime.h`, which lives in the `cuda-cudart` component. An +earlier revision of this rule left it out, and nvcc resolved `cuda_runtime.h` +from `/usr/include` and then read the host's `crt/host_config.h` beside it. The +build failed with the host toolkit's complaint while using the payload's +compiler. + +## Two routes, and why the primary one is clang + +The rule package compiles the device unit either way: + +- **clang** (`-x cuda --cuda-path=`) is the default and what + `[toolchain] default = "llvm@22.1.8"` selects. The compiler that builds the + rest of the project builds the device unit too: no second host compiler and + no host-compiler bound. + + It does pass one flag of NVIDIA's own. A device unit that includes + `` stops at `crt/host_defines.h:67` with `"libc++ is not + supported on x86 system"`, because that guard reads `__CUDACC__` — which + clang defines when it compiles CUDA — and an LLVM toolchain's clang uses + libc++. The refusal is about nvcc's host pass, so the rule passes + `-D_ALLOW_UNSUPPORTED_LIBCPP` on this route only. This example's own kernel + never showed it: a bare kernel includes no toolkit header at all. +- **nvcc** (`-ccbin `) is taken when the project's toolchain is GCC. + It drives a second compiler, and that is where its constraints come from. + +`MCPP_EXAMPLE_CUDA_ROUTE=clang|nvcc` overrides the choice, and the rule +declares `rerun_if_env_changed` for it. + +Two pairings nvcc cannot have, both stated before the compile rather than +discovered inside it: + +- **A host compiler past the bound.** nvcc states a maximum GCC major in its + own `crt/host_config.h`. The rule reads it, uses the project's toolchain when + it fits, otherwise a `xim:gcc` payload the project declared for this purpose, + and otherwise refuses naming the declaration to add. Measured: GCC 16 under + nvcc 12.9 fails inside GCC's own `` even with + `-allow-unsupported-compiler` — that escape hatch admits a compiler one step + past the bound, not a standard library two majors newer. +- **An old toolkit and a new C library.** Toolkit 12.9's + `crt/math_functions.h` redeclares the C23 functions `cospi`, `sinpi` and + `rsqrt` for the host without `noexcept`; glibc 2.41 and later declare them + with it, and since C++17 that is part of the function type. The compile stops + with six `exception specification is incompatible` errors naming a glibc + header and a CUDA header, and no decision. The rule reads the C library's + `bits/mathcalls.h` through `mcpp::toolchain_sysroot()` and refuses the pair, + naming the 13.x toolkit as the way out. The clang route does not include that + header at all. + +A second compiler also has to be told where it is. `mcpp::toolchain_sysroot()` +and `mcpp::toolchain_binutils_dir()` are the `--sysroot` and `-B` mcpp passes +to its own compiler; without forwarding them, NVIDIA's `crt/host_config.h` +stops at `features.h: No such file or directory`. + +Everything about a compiler's spelling lives in the rule package. The engine +owns the graph, the artifact's identity and the architecture set; it does not +own `-gencode` or `--cuda-gpu-arch`. + +## What the rule reports before the first compile + +The rule states machine facts through the build program's own channel, and mcpp +compares them: -Everything about nvcc's spelling lives in the rule package. The engine owns the -graph, the artifact's identity and the architecture set; it does not own -`-gencode`. +``` +mcpp:fact=cuda.driver=12.4 +mcpp:floor=cuda.driver >= 12.0 +``` + +The fact comes from opening the driver's own library through the +`libcuda-host-link` sentinel and asking it for its version; the floor comes +from the toolkit the project named. mcpp refuses a build whose floor is not met +and says so in one sentence, because the failure it prevents is not a build +failure: + +``` +error: `cuda-saxpy` requires cuda.driver >= 13.0, and this machine has 12.4. +``` + +A separate advisory covers PTX: embedded PTX emitted by a toolkit newer than +the driver cannot be JIT-compiled by that driver, so hardware outside the named +architecture set will not run. The named architectures still do, so this is a +warning rather than a refusal. + +The engine holds no vendor name for any of this. `cuda.driver` is a string +flowing from a declaration to a comparison; a second backend needs no engine +change. A unit test refuses vendor probes in `src/`. ## Verified -On an NVIDIA RTX 4080 (compute capability 8.9) with CUDA 12.0 and driver -550.144.03: +On an NVIDIA RTX 4080 (compute capability 8.9), driver 550.144.03 reporting +CUDA 12.4, with an LLVM 22.1.8 toolchain: ``` $ mcpp run - Running `target/.../bin/cuda-saxpy` + Running `target/x86_64-linux-gnu//bin/cuda-saxpy` +12 24 36 48 + +$ mcpp run --no-accel + Running `target/x86_64-linux-gnu//bin/cuda-saxpy` 12 24 36 48 ``` -which is `2.0 * [1,2,3,4] + [10,20,30,40]` computed on the device. +which is `2.0 * [1,2,3,4] + [10,20,30,40]`, computed on the device in the first +case and on the host in the second. The two artifact directories differ, and +the CPU one contains no `cudaMalloc`. + +The nvcc route is not exercisable on that machine: the 12.9 toolkit meets the +driver and not the C library, and the 13.3 toolkit meets the C library and not +the driver. Both refusals are the ones described above, and both name the way +out. ## Where the driver comes from diff --git a/examples/09-cuda-kernel/app/build.mcpp b/examples/09-cuda-kernel/app/build.mcpp index d263f4841..7a1b989c8 100644 --- a/examples/09-cuda-kernel/app/build.mcpp +++ b/examples/09-cuda-kernel/app/build.mcpp @@ -1,20 +1,21 @@ import std; import mcpp; -import example.rules.cuda; +import mcpplibs.rules.cuda; +// Everything the rule needs is in the manifest: the architectures in +// `[build] accel`, the device sources in the constrained glob, the toolkit +// under `[xlings.workspace]`. This program names the island's include +// directory and says "go". +// +// The route follows the project's toolchain (clang for an LLVM toolchain, +// nvcc for a GCC one); MCPP_EXAMPLE_CUDA_ROUTE=nvcc|clang overrides it, which +// is how the alternate route is measured without editing the manifest. int main() { - mcpp::rerun_if_changed_glob("src/kernels/**/*.cu"); - mcpp::rerun_if_changed_glob("include/**/*.h"); - - example::rules::cuda::options opt; - // RTX 4080 is compute capability 8.9. Named rather than detected: what a - // build compiles for is a decision, and a machine's own hardware is a poor - // default for it — the artifact would run here and nowhere else. - opt.archs = { "sm_89" }; - // Embed the portable form as well, so the same object runs on hardware - // newer than this one. - opt.ptx = "89"; + mcpp::rerun_if_env_changed("MCPP_EXAMPLE_CUDA_ROUTE"); + mcpplibs::rules::cuda::options opt; opt.includes = { "include" }; - const std::vector sources{ "src/kernels/saxpy.cu" }; - return example::rules::cuda::compile(sources, opt) ? 0 : 1; + if (const char* r = std::getenv("MCPP_EXAMPLE_CUDA_ROUTE"); r && *r) + opt.which = std::string_view(r) == "nvcc" ? mcpplibs::rules::cuda::route::nvcc + : mcpplibs::rules::cuda::route::clang; + return mcpplibs::rules::cuda::compile(opt) ? 0 : 1; } diff --git a/examples/09-cuda-kernel/app/mcpp.toml b/examples/09-cuda-kernel/app/mcpp.toml index 8ca590cfc..5118ffab5 100644 --- a/examples/09-cuda-kernel/app/mcpp.toml +++ b/examples/09-cuda-kernel/app/mcpp.toml @@ -2,7 +2,7 @@ name = "cuda-saxpy" namespace = "example" version = "0.1.0" -description = "A CUDA kernel behind a seam module" +description = "A CUDA kernel behind a seam module, with a CPU fallback" accelerators = ["cuda"] [language] @@ -10,39 +10,63 @@ standard = "c++23" modules = true import_std = true +# The primary route is clang: the toolchain's own clang++ compiles the device +# unit (`-x cuda`), so there is no second host compiler and no host-compiler +# bound to satisfy. With a GCC toolchain the rule package takes the nvcc route +# instead, driving that GCC and reading the bound nvcc states for it. +[toolchain] +default = "llvm@22.1.8" + [dependencies] rules-cuda = { path = "../rules-cuda", host-module = true } # The driver's userspace library, reached through an index package that owns # the one hop mcpp needs: a directory on the artifact's runtime search path. # mcpp's private loader does not consult /usr/lib, so without it the statically -# linked CUDA runtime cannot dlopen the driver. +# linked CUDA runtime cannot dlopen the driver. It is the one CUDA component +# that cannot be an ordinary payload -- the licence forbids redistributing it +# and it is in ABI lockstep with the kernel module. [dependencies.compat] cuda-runtime = "2026.09.05" -# The driver's userspace library, reached through the sentinel package. -# -# It is the one CUDA component that cannot be an ordinary payload: NVIDIA's -# driver licence forbids redistributing it, and it is in ABI lockstep with the -# kernel module, so a version of it is meaningless outside the machine it came -# from. The sentinel installs a symlink to whatever the host has, which gives -# 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. +# The toolkit this project builds with, named rather than discovered. These +# are PAYLOADS, so the version is the project's choice and not the machine's. +# The 12.9 line is named on purpose: a runtime must not be newer than the +# driver it will meet, and 12.x reaches every driver from r525 onward. The +# rule package states the driver relation and mcpp compares it before the +# first compile. [xlings.workspace] +"xim:cuda-nvcc" = "12.9.86" +"xim:cuda-cudart" = "12.9.79" "xim:libcuda-host-link" = { linux = "0.0.1" } [build] -# The CUDA runtime is linked STATICALLY. mcpp refuses a dynamic link against -# the host's libcudart because its private loader does not consult /usr/lib, -# and it is right to: such an artifact is not self-contained. Linking the -# redistributable half in leaves exactly one host dependency, libcuda.so.1, -# which is the driver and genuinely cannot be redistributed — that is what the -# libcuda-host-link sentinel package in xim exists for. -accel = "cuda12.0+{sm_89} ptx>=89" +# What this build compiles device code FOR. Written once, here: the rule +# package derives its own flags (`--cuda-gpu-arch`, `-gencode`) from it. +accel = "cuda12.9+{sm_89} ptx>=89" +# The device sources carry the accel they are for. Under `--no-accel` the glob +# is left out and the CPU fallback below takes its place; under an accel that +# does not cover it the build is refused naming both. +sources = [ + "src/*.cppm", + "src/*.cpp", + { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }, +] include_dirs = ["include"] -ldflags = ["-L/usr/local/cuda/lib64", "-L/usr/lib/x86_64-linux-gnu", - "-lcudart_static", "-lrt", "-lpthread", "-ldl"] + +# The CUDA runtime is linked STATICALLY, and only when a device build asks for +# it. Linking the redistributable half in leaves exactly one host dependency, +# libcuda.so.1, which the driver owns and the sentinel package reaches. +# ⭐ NO ABSOLUTE PATHS: the rule package puts the payload's library directory on +# the link line from `mcpp::xpkg_dir`, so this manifest names libraries only. +[target.'cfg(accelerator = "cuda")'.build] +ldflags = ["-lcudart_static", "-lrt", "-lpthread", "-ldl"] + +# The CPU-only variant: the same seam, a host implementation behind it. This +# is what `mcpp build --no-accel` produces, and what a machine with no device +# runs. +[target.'cfg(not(accelerator = "cuda"))'.build] +sources = ["src/cpu/*.cpp"] [targets.cuda-saxpy] kind = "bin" diff --git a/examples/09-cuda-kernel/app/src/cpu/saxpy.cpp b/examples/09-cuda-kernel/app/src/cpu/saxpy.cpp new file mode 100644 index 000000000..4de4ad625 --- /dev/null +++ b/examples/09-cuda-kernel/app/src/cpu/saxpy.cpp @@ -0,0 +1,11 @@ +// The CPU implementation behind the same seam. Compiled only when the build +// asks for no accelerator (`mcpp build --no-accel`), through the +// `cfg(not(accelerator = "cuda"))` section of the manifest; the device island +// and this file define the same symbol and are never in one link. +#include "saxpy/saxpy.h" + +extern "C" int saxpy_device(float a, const float* x, const float* y, + float* out, unsigned n) { + for (unsigned i = 0; i < n; ++i) out[i] = a * x[i] + y[i]; + return 0; +} diff --git a/examples/09-cuda-kernel/rules-cuda/mcpp.toml b/examples/09-cuda-kernel/rules-cuda/mcpp.toml index cdfa36abf..eb6012051 100644 --- a/examples/09-cuda-kernel/rules-cuda/mcpp.toml +++ b/examples/09-cuda-kernel/rules-cuda/mcpp.toml @@ -1,6 +1,6 @@ [package] name = "rules-cuda" -namespace = "example" +namespace = "mcpplibs" version = "0.1.0" -description = "Compile CUDA device translation units with nvcc (role = object)" +description = "Compile CUDA device translation units: clang -x cuda by default, nvcc as the alternate route (role = object)" license = "Apache-2.0" diff --git a/examples/09-cuda-kernel/rules-cuda/src/rules-cuda.cppm b/examples/09-cuda-kernel/rules-cuda/src/rules-cuda.cppm index da1f460f8..6b537d7cf 100644 --- a/examples/09-cuda-kernel/rules-cuda/src/rules-cuda.cppm +++ b/examples/09-cuda-kernel/rules-cuda/src/rules-cuda.cppm @@ -2,77 +2,275 @@ // // WHY A RULE PACKAGE RATHER THAN THE ENGINE // -// Everything below is knowledge about one vendor's driver: where nvcc lives, -// how it spells an architecture, which host compilers it tolerates, and how it -// must be told about them. None of it is knowledge about the build graph. The -// engine owns the graph, the artifact's identity and the architecture set; the -// spelling of the command that produces the object is this file's business. +// Everything below is knowledge about one vendor's tools: where the toolkit's +// pieces live, how an architecture is spelled, which host compilers nvcc +// tolerates, how a driver states its version. None of it is knowledge about +// the build graph. The engine owns the graph, the artifact's identity and the +// architecture set (`[build] accel`); the spelling of the command that +// produces an object, and every probe of the machine, is this file's business. +// `tests/unit/test_core_vendor_probes.cpp` in mcpp holds that line from the +// other side: the engine names no vendor tool. // -// The division is not aesthetic. CMake carries nine years of open issues about -// -Xcompiler escaping, response files and device-link edge cases, and every one -// of them is a fact about nvcc that ended up inside a build system. +// TWO ROUTES, ONE PRIMARY // -// THE HOST COMPILER IS THE PART THAT SURPRISES PEOPLE +// `clang -x cuda` is the primary route: the resolved toolchain's own clang +// compiles the device unit, so there is no second host compiler and no +// host-compiler bound to satisfy. nvcc is the alternate, taken when the +// project's toolchain is GCC or when asked for: it drives a host compiler +// (`-ccbin`) and refuses one newer than the bound its `crt/host_config.h` +// states, which this rule reads and reports. // -// nvcc refuses host compilers newer than a bound it states in its own -// crt/host_config.h, and mcpp's toolchain payload is routinely newer than that -// bound. Passing mcpp's compiler through would fail; passing nothing would let -// nvcc pick the system default and fail the same way. So the rule reads the -// bound and selects a host compiler that satisfies it, and says which one it -// chose — an unexplained third compiler in a build is worse than an error. +// WHAT THE RULE TELLS THE ENGINE +// +// Objects, through `mcpp::action` (role "object"); library directories, +// through `mcpp::link_search`; and three claims about the machine that the +// engine compares or relays before the first compile: +// - the driver's version, read through the driver's own library and stated +// with `mcpp::fact`, together with the floor the runtime needs +// (`mcpp::floor`); an unmet floor refuses the build with both values; +// - whether nvcc can reach its own back-end stages (`--dryrun`), as an +// advisory naming the first stage that does not resolve; +// - whether the embedded PTX can be JIT-compiled by this driver, as an +// advisory, because the SASS for the named architectures still runs. + +module; +#include +#include +#if !defined(_WIN32) +#include +#endif -export module example.rules.cuda; +export module mcpplibs.rules.cuda; import std; import mcpp; -export namespace example::rules::cuda { +export namespace mcpplibs::rules::cuda { + +enum class route { automatic, clang, nvcc }; struct options { - // Device architectures to emit real code for. No default: the set a build - // compiles for is not the set the machine happens to have, and guessing - // one produces an artifact that runs here and nowhere else. - std::vector archs; - // The virtual architecture to embed a portable form of, so newer hardware - // can JIT. Empty means none, and then the artifact runs only on `archs`. - std::string ptx; - // Header search paths for the island, relative to the package root. The - // island's own interface lives in one of these, and nvcc is a separate - // driver that inherits nothing from the C++ side's include configuration. + route which = route::automatic; + // Header search paths for the island. Relative entries resolve against the + // package root; an ABSOLUTE entry is passed through unchanged. + // + // ⭐ THE ABSOLUTE FORM IS FOR A DEPENDENCY'S HEADERS. A device compiler is + // a separate driver and inherits nothing from the C++ side's include + // configuration, so a package whose device code includes a dependency's + // header -- ggml's CUDA backend includes `cublas_v2.h` -- has to name that + // dependency's directory here, and it knows it only as the absolute path + // `mcpp::dep_dir` answered with. std::vector includes; - std::string out_dir = std::string(mcpp::out_dir()); + std::string out_dir = std::string(mcpp::out_dir()); }; -struct edge { - std::string id, description; - std::vector command, inputs, outputs; +// ─── What the engine said ────────────────────────────────────────────────── + +// The `cuda` chunk of `mcpp::accel()`, in this rule's own reading: the engine +// carries the string and compares it as a shape; what `sm_89` means is ours. +struct target { + std::string version; // "12.9" + std::vector archs; // {"sm_89"} + std::string ptx; // "89" when a portable form is embedded + bool present = false; }; +// Split on one character. Written out rather than taken from : GCC 16 +// refuses the ranges split view instantiated inside an exported inline function +// when build.mcpp imports this module (`conflicting deduced return type for +// imported declaration ... view_interface::data()`), and clang does not. +inline std::vector split(std::string_view s, char sep) { + std::vector out; + for (std::size_t i = 0; i <= s.size();) { + auto j = s.find(sep, i); + out.push_back(s.substr(i, j == std::string_view::npos ? s.size() - i : j - i)); + if (j == std::string_view::npos) break; + i = j + 1; + } + return out; +} + +inline std::string_view trim(std::string_view s) { + while (!s.empty() && (s.front() == ' ' || s.front() == '\t')) s.remove_prefix(1); + while (!s.empty() && (s.back() == ' ' || s.back() == '\t')) s.remove_suffix(1); + return s; +} + +inline target parse_target(std::string_view accel) { + target t; + for (std::size_t i = 0; i <= accel.size();) { + auto comma = accel.find(',', i); + auto open = accel.find('{', i), close = accel.find('}', i); + if (open != std::string_view::npos && close != std::string_view::npos + && comma != std::string_view::npos && comma > open && comma < close) + comma = accel.find(',', close); + auto chunk = trim(comma == std::string_view::npos ? accel.substr(i) + : accel.substr(i, comma - i)); + i = comma == std::string_view::npos ? accel.size() + 1 : comma + 1; + if (!chunk.starts_with("cuda")) continue; + t.present = true; + auto plus = chunk.find('+'); + t.version = std::string(trim(chunk.substr(4, plus == std::string_view::npos + ? chunk.size() - 4 : plus - 4))); + if (plus != std::string_view::npos) { + auto o = chunk.find('{', plus), c = chunk.find('}', plus); + if (o != std::string_view::npos && c != std::string_view::npos) + for (auto part : split(chunk.substr(o + 1, c - o - 1), ',')) + if (auto a = trim(part); !a.empty()) t.archs.emplace_back(a); + auto tail = chunk.substr(c == std::string_view::npos ? chunk.size() : c + 1); + for (auto key : {"ptx>=", "floor>="}) + if (auto p = tail.find(key); p != std::string_view::npos) + t.ptx = std::string(trim(tail.substr(p + std::string_view(key).size()))); + } + } + return t; +} + +inline std::vector device_sources() { + std::vector out; + for (auto part : split(std::string_view(mcpp::device_sources()), '\n')) + if (auto s = trim(part); !s.empty()) out.emplace_back(s); + return out; +} + // ─── Locating the toolkit ────────────────────────────────────────────────── -inline std::string first_existing(std::span candidates) { - for (auto const& c : candidates) - if (!c.empty() && std::filesystem::exists(c)) return c; +// The toolkit this project declared under `[xlings.workspace]`, by component. +// The 13.x line splits the compiler across `cuda-nvcc`, `cuda-crt` and +// `libnvvm`; the 12.x line keeps them in `cuda-nvcc`. Either way the project +// names the compiler and this rule finds the pieces. +struct toolkit { + std::string nvcc_root, cudart_root, crt_root, driver_dir; + std::string nvcc() const { return nvcc_root + "/bin/nvcc"; } + std::string host_config() const { + for (auto const* r : { &crt_root, &nvcc_root, &cudart_root }) { + if (r->empty()) continue; + auto p = *r + "/include/crt/host_config.h"; + if (std::filesystem::exists(p)) return p; + } + return {}; + } + std::vector include_dirs() const { + std::vector out; + for (auto const* r : { &cudart_root, &crt_root, &nvcc_root }) + if (!r->empty() && std::filesystem::is_directory(*r + "/include")) + out.push_back(*r + "/include"); + return out; + } + std::vector lib_dirs() const { + std::vector out; + for (auto const* r : { &cudart_root, &nvcc_root }) + for (auto const* sub : { "/lib", "/lib64" }) + if (!r->empty() && std::filesystem::is_directory(*r + sub)) + out.push_back(*r + sub); + return out; + } +}; + +inline std::string xpkg(const char* name) { + if (const char* d = mcpp::xpkg_dir("xim", name); d && *d) return d; return {}; } -inline std::string find_nvcc() { - std::vector c; - for (const char* var : { "CUDA_PATH", "CUDA_HOME" }) - if (const char* v = std::getenv(var)) c.push_back(std::string(v) + "/bin/nvcc"); - c.push_back("/usr/local/cuda/bin/nvcc"); - c.push_back("/usr/bin/nvcc"); - return first_existing(c); +inline std::optional find_toolkit() { + toolkit t; + t.nvcc_root = xpkg("cuda-nvcc"); + t.cudart_root = xpkg("cuda-cudart"); + t.crt_root = xpkg("cuda-crt"); + t.driver_dir = xpkg("libcuda-host-link"); + if (t.nvcc_root.empty() || t.cudart_root.empty()) { + std::println(std::cerr, + "mcpplibs.rules.cuda: the toolkit is not declared.\n" + " Name it under [xlings.workspace] and mcpp provisions it on first use:\n" + " \"xim:cuda-nvcc\" = \"12.9.86\"\n" + " \"xim:cuda-cudart\" = \"12.9.79\"\n" + " (found nvcc: '{}', cudart: '{}')", t.nvcc_root, t.cudart_root); + return std::nullopt; + } + return t; } -inline std::string find_host_config(std::string_view nvcc) { - std::vector c; - if (!nvcc.empty()) { - std::filesystem::path p{std::string(nvcc)}; - c.push_back((p.parent_path().parent_path() / "include/crt/host_config.h").string()); +// ─── Probes: what the machine has, what the toolkit needs ────────────────── + +// The driver's version through the driver's own library, reached through the +// sentinel package rather than /usr/lib. "" when there is no driver here, +// which is a fact about the machine and not a failure of the build. +inline std::string driver_version(const toolkit& t) { +#if defined(_WIN32) + return {}; +#else + if (t.driver_dir.empty()) return {}; + const auto lib = t.driver_dir + "/lib/libcuda.so.1"; + if (!std::filesystem::exists(lib)) return {}; + // What would change the answer is the library the answer was read from. + mcpp::rerun_if_changed(lib.c_str()); + void* h = ::dlopen(lib.c_str(), RTLD_LAZY | RTLD_LOCAL); + if (!h) { + const char* why = ::dlerror(); + mcpp::warning(std::format("could not open the driver library {}: {}", lib, + why ? why : "(no reason given)").c_str()); + return {}; } - c.push_back("/usr/include/crt/host_config.h"); - return first_existing(c); + using fn = int (*)(int*); + auto get = reinterpret_cast(::dlsym(h, "cuDriverGetVersion")); + int v = 0; + std::string out; + if (get && get(&v) == 0 && v > 0) out = std::format("{}.{}", v / 1000, (v % 1000) / 10); + ::dlclose(h); + return out; +#endif +} + +inline int major_of(std::string_view v) { + int m = 0; + for (char c : v) { if (!std::isdigit(static_cast(c))) break; m = m * 10 + (c - '0'); } + return m; +} + +inline bool version_at_least(std::string_view have, std::string_view want) { + auto parse = [](std::string_view s) { + std::vector out; int acc = 0; bool digits = false; + for (char c : s) { + if (c == '.') { out.push_back(acc); acc = 0; digits = false; continue; } + if (!std::isdigit(static_cast(c))) break; + acc = acc * 10 + (c - '0'); digits = true; + } + if (digits) out.push_back(acc); + return out; + }; + auto h = parse(have), w = parse(want); + for (std::size_t i = 0; i < std::max(h.size(), w.size()); ++i) { + int a = i < h.size() ? h[i] : 0, b = i < w.size() ? w[i] : 0; + if (a != b) return a > b; + } + return true; +} + +// State the driver relation. The engine compares the floor against the fact +// and refuses with both values; this rule only knows which numbers matter. +// +// The floor is the toolkit's major: a 12.x runtime runs on any 12.x driver +// (minor-version compatibility), and fails at the first allocation on an 11.x +// one. The embedded PTX is a separate, softer question: PTX emitted by toolkit +// 12.9 is JIT-compiled only by a driver at or above 12.9, but the SASS for the +// named architectures still runs, so a driver below the toolkit costs reach on +// newer hardware rather than correctness here -- reported, not enforced. +inline void state_driver_relation(const toolkit& t, const target& tg) { + const auto driver = driver_version(t); + if (!driver.empty()) mcpp::fact("cuda.driver", driver.c_str()); + else mcpp::warning("no driver library reachable through xim:libcuda-host-link; " + "the build proceeds and the artifact will find no device at run time"); + const int major = major_of(tg.version); + if (major > 0) mcpp::floor(std::format("cuda.driver >= {}.0", major).c_str()); + if (!driver.empty() && !tg.ptx.empty() && !version_at_least(driver, tg.version)) + mcpp::warning(std::format( + "the PTX embedded for compute_{} was emitted by toolkit {} and this driver " + "serves {}; hardware newer than {{{}}} will not be able to JIT it. The named " + "architectures run. Build with a toolkit at or below the driver, or add the " + "newer hardware's SASS to [build] accel.", + tg.ptx, tg.version, driver, [&] { + std::string s; for (auto& a : tg.archs) { if (!s.empty()) s += ','; s += a; } + return s; }()).c_str()); } // The greatest gcc major and the greatest clang major the toolkit accepts. @@ -95,86 +293,291 @@ inline bounds read_bounds(std::string_view headerPath) { } return n ? v : 0; }; - if (auto p = text.find("__GNUC__ > "); p != std::string::npos) - b.gcc = number_after(p + 10); + if (auto p = text.find("__GNUC__ > "); p != std::string::npos) b.gcc = number_after(p + 10); if (auto p = text.find("clang version must be less than "); p != std::string::npos) if (int excl = number_after(p + 31); excl > 0) b.clang = excl - 1; return b; } -// A host compiler this toolkit accepts, or empty. +// Does the C library this build compiles against declare the C23 functions +// `cospi`, `sinpi` and `rsqrt`? +// +// ⚠️ Measured 2026-09-05 against glibc 2.44. Toolkit 12.9's +// `crt/math_functions.h` declares those same names for the host WITHOUT +// `noexcept`; glibc declares them WITH it, and since C++17 that is part of the +// function type. nvcc's front end stops with six `exception specification is +// incompatible` errors that name a glibc header and a CUDA header and leave +// the reader to work out that neither is at fault alone. The 13.x line does +// not redeclare them and compiles cleanly against the same C library. // -// Searched newest-first within the bound, because a newer accepted compiler -// produces better host code than an older one and both are equally correct. -inline std::string pick_host_compiler(const bounds& b) { - for (int v = b.clang ? b.clang : 20; v >= 3; --v) { - auto p = std::format("/usr/bin/clang++-{}", v); - if ((b.clang == 0 || v <= b.clang) && std::filesystem::exists(p)) return p; +// Read, not probed. The answer is one substring of one header the sysroot +// already contains; a probe compile would spend a second nvcc invocation to +// learn the same thing, and would report it as a compile failure rather than +// as a pairing that cannot work. +inline bool libc_declares_c23_pi_math(std::string_view sysroot) { + if (sysroot.empty()) return false; + for (auto const* rel : { "/usr/include/bits/mathcalls.h", "/include/bits/mathcalls.h" }) { + std::ifstream in{std::string(sysroot) + rel}; + if (!in) continue; + std::string text{std::istreambuf_iterator(in), std::istreambuf_iterator()}; + return text.find("(cospi,") != std::string::npos + && text.find("(rsqrt,") != std::string::npos; } - for (int v = b.gcc ? b.gcc : 20; v >= 5; --v) { - auto p = std::format("/usr/bin/g++-{}", v); - if ((b.gcc == 0 || v <= b.gcc) && std::filesystem::exists(p)) return p; + return false; +} + +// ` -dumpversion` → major. The host compiler is the toolchain mcpp +// resolved for this build, so its version is a fact of the build, not a guess. +inline int compiler_major(const std::string& cc) { +#if defined(_WIN32) + (void)cc; return 0; +#else + std::string cmd = cc + " -dumpversion 2>/dev/null"; + if (FILE* p = ::popen(cmd.c_str(), "r")) { + char buf[64] = {}; + std::string s; + if (std::fgets(buf, sizeof buf, p)) s = buf; + ::pclose(p); + return major_of(s); } - return {}; + return 0; +#endif +} + +// The first back-end stage nvcc names but cannot resolve, from its own plan. +// nvcc invokes cicc, cudafe++, ptxas and fatbinary by bare name on a PATH it +// states in the plan; a stage that does not resolve there fails the compile +// with `sh: 1: cicc: not found`, naming nothing that helps. +inline std::optional unreachable_stage(const toolkit& t, const std::string& ccbin) { +#if defined(_WIN32) + (void)t; (void)ccbin; return std::nullopt; +#else + const auto probe = std::filesystem::temp_directory_path() / "mcpp-rules-cuda-dryrun.cu"; + { std::ofstream(probe) << "__global__ void k() {}\n"; } + std::string cmd = std::format("{} --dryrun -ccbin {} -c {} -o /dev/null 2>&1", + t.nvcc(), ccbin, probe.string()); + std::string text; + if (FILE* p = ::popen(cmd.c_str(), "r")) { + char buf[4096]; + while (std::fgets(buf, sizeof buf, p)) text += buf; + ::pclose(p); + } + std::filesystem::remove(probe); + std::string path; + std::vector stages; + for (auto l : split(text, '\n')) { + if (!l.starts_with("#$ ")) continue; + l.remove_prefix(3); + if (l.starts_with("PATH=")) { path = std::string(l.substr(5)); continue; } + for (auto const* stage : { "cicc", "cudafe++", "ptxas", "fatbinary", "nvlink" }) { + auto pos = l.find(stage); + if (pos == 0 || (pos != std::string_view::npos && (l[pos - 1] == ' ' || l[pos - 1] == '"'))) + if (std::ranges::find(stages, stage) == stages.end()) stages.emplace_back(stage); + } + } + if (stages.empty()) return std::nullopt; // no plan, no finding + for (auto const& stage : stages) { + bool found = false; + for (auto dir : split(path, ':')) { + std::string d(dir); + if (!d.empty() && std::filesystem::exists(d + "/" + stage)) { found = true; break; } + } + if (!found) return stage; + } + return std::nullopt; +#endif } // ─── Planning ────────────────────────────────────────────────────────────── +struct edge { + std::string id, description; + std::vector command, inputs, outputs; +}; + +inline route decide(route asked) { + if (asked != route::automatic) return asked; + return std::string_view(mcpp::compiler()) == "clang" ? route::clang : route::nvcc; +} + inline std::vector plan(std::span sources, options opt = {}) { std::vector out; const std::string root = mcpp::manifest_dir(); if (root.empty()) { - std::println(std::cerr, "example.rules.cuda: no mcpp build context — " - "this runs from build.mcpp"); - return out; - } - if (opt.archs.empty()) { - std::println(std::cerr, - "example.rules.cuda: no architectures given.\n" - " The set a build compiles for is not the set this machine has, so\n" - " there is no default. Pass e.g. archs = {{\"sm_89\"}}."); + std::println(std::cerr, "mcpplibs.rules.cuda: no mcpp build context -- this runs from build.mcpp"); return out; } - const auto nvcc = find_nvcc(); - if (nvcc.empty()) { + const auto tg = parse_target(mcpp::accel()); + if (!tg.present || tg.archs.empty()) { + // C19: a device build that names no device is refused HERE, not at + // run time as `no kernel image is available for execution`. std::println(std::cerr, - "example.rules.cuda: no nvcc found (looked at $CUDA_PATH/bin, " - "/usr/local/cuda/bin, /usr/bin)."); + "mcpplibs.rules.cuda: [build] accel names no CUDA architecture (accel = \"{}\").\n" + " Write e.g. accel = \"cuda12.9+{{sm_89}} ptx>=89\" -- the set a build compiles\n" + " for is a decision, and the machine's own hardware is a poor default for it.", + mcpp::accel()); return out; } - const auto b = read_bounds(find_host_config(nvcc)); - const auto ccbin = pick_host_compiler(b); - if (ccbin.empty()) { - std::println(std::cerr, - "example.rules.cuda: this toolkit accepts gcc <= {} and clang <= {}, " - "and no such host compiler is installed.", b.gcc, b.clang); - return out; + auto tk = find_toolkit(); + if (!tk) return out; + state_driver_relation(*tk, tg); + + const route r = decide(opt.which); + const std::string tcdir = mcpp::toolchain_dir(); + std::string driver_cc; // the compiler that runs the device unit + std::vector front; // the command up to the input file + if (r == route::clang) { + driver_cc = tcdir + "/bin/clang++"; + if (!std::filesystem::exists(driver_cc)) { + std::println(std::cerr, "mcpplibs.rules.cuda: the clang route needs the toolchain's clang++ at {}", driver_cc); + return out; + } + front = { driver_cc, "-x", "cuda", "-std=c++17", "-O2", "-fPIC", + "--cuda-path=" + tk->nvcc_root, "-Wno-unknown-cuda-version", + // ⚠️ NVIDIA'S HEADER REFUSES libc++, AND THE REFUSAL IS + // ABOUT nvcc RATHER THAN ABOUT THIS COMPILER. + // + // crt/host_defines.h:67: error: "libc++ is not supported + // on x86 system" + // + // The guard is `#if defined(__CUDACC__) && … && + // defined(_LIBCPP_VERSION)`, and clang defines `__CUDACC__` + // when it compiles CUDA itself — so a device unit that + // includes stops here on any LLVM + // toolchain, which is the toolchain this route exists for. + // Measured on ggml's CUDA backend; the CUDA example's own + // kernel never showed it because a bare kernel includes no + // toolkit header at all. + // + // The escape hatch is upstream's own, and it is passed only + // on this route: nvcc's host pass really does break against + // libc++, and nothing here weakens that. + "-D_ALLOW_UNSUPPORTED_LIBCPP" }; + for (auto const& inc : tk->include_dirs()) front.push_back("-I" + inc); + for (auto const& a : tg.archs) front.push_back("--cuda-gpu-arch=" + a); + // clang checks ptxas and fatbinary itself; say so before it does. + for (auto const* tool : { "ptxas", "fatbinary" }) + if (!std::filesystem::exists(tk->nvcc_root + "/bin/" + tool)) + mcpp::warning(std::format("the toolkit payload has no {}; clang invokes it " + "after generating PTX", tool).c_str()); + std::println("mcpplibs.rules.cuda: clang route -- {} (toolkit {})", driver_cc, tk->nvcc_root); + } else { + // nvcc drives the toolchain's own compiler, and refuses one newer than + // the bound its header states. Read the bound; if exceeded, pass the + // escape hatch and say so -- an unexplained flag is worse than a note. + const bool clangHost = std::string_view(mcpp::compiler()) == "clang"; + if (clangHost) { + // Measured: nvcc's own crt/host_defines.h stops the compile with + // `libc++ is not supported on x86 system`, and libc++ is what an + // LLVM toolchain's clang uses. The pairing that works is nvcc with + // a GCC toolchain; with an LLVM toolchain the clang route is the + // one to take, and it is the default. + std::println(std::cerr, + "mcpplibs.rules.cuda: the nvcc route needs a GCC host compiler; this project's " + "toolchain is LLVM, whose clang uses libc++ and nvcc refuses it. Use the clang " + "route (the default for an LLVM toolchain) or set [toolchain] to a gcc payload."); + return out; + } + // The other pairing this route cannot have: an old toolkit and a C + // library new enough to have the C23 `pi` functions. Stated before the + // compile, because the compile's own report names two headers and no + // decision. + if (major_of(tg.version) < 13 + && libc_declares_c23_pi_math(mcpp::toolchain_sysroot())) { + std::println(std::cerr, + "mcpplibs.rules.cuda: toolkit {} redeclares the C23 functions cospi, sinpi and " + "rsqrt for the host without `noexcept`, and the C library this build compiles " + "against declares them with it; nvcc's front end refuses the pair.\n" + " Name a 13.x toolkit, whose headers leave them to the C library:\n" + " [xlings.workspace]\n" + " \"xim:cuda-nvcc\" = \"13.3.33\"\n" + " \"xim:cuda-crt\" = \"13.3.33\"\n" + " \"xim:cuda-cudart\" = \"13.3.29\"\n" + " or take the clang route, which does not include that header at all.", + tg.version); + return out; + } + // The host compiler nvcc drives, chosen within the bound the toolkit + // states. Measured: gcc 16 under nvcc 12.9 (bound gcc <= 14) fails inside + // nvcc's front end on GCC 16's own even with + // -allow-unsupported-compiler -- the escape hatch admits a compiler one + // step past the bound, not a standard library two majors newer. So the + // rule does not guess: the toolchain's g++ when it is within the bound, + // otherwise a gcc payload the project declared for this purpose, and + // otherwise a refusal that says which declaration to add. + const auto b = read_bounds(tk->host_config()); + const std::string tcGcc = tcdir + "/bin/g++"; + const int tcMajor = compiler_major(tcGcc); + if (b.gcc == 0 || tcMajor <= b.gcc) { + driver_cc = tcGcc; + } else if (auto payload = xpkg("gcc"); !payload.empty() + && compiler_major(payload + "/bin/g++") <= b.gcc) { + driver_cc = payload + "/bin/g++"; + mcpp::warning(std::format( + "nvcc {} states gcc <= {} in {}; the toolchain's gcc {} exceeds it, so the " + "device unit is compiled with the declared xim:gcc payload ({}). The clang " + "route has no such bound.", tg.version, b.gcc, tk->host_config(), tcMajor, + driver_cc).c_str()); + } else { + std::println(std::cerr, + "mcpplibs.rules.cuda: nvcc {} accepts gcc <= {} ({}), and this project's " + "toolchain is gcc {}.\n" + " Declare a gcc payload within the bound and the rule drives that one:\n" + " [xlings.workspace]\n" + " \"xim:gcc\" = \"13.3.0\"\n" + " or take the clang route with [toolchain] default = \"llvm@22.1.8\".", + tg.version, b.gcc, tk->host_config(), tcMajor); + return out; + } + front = { tk->nvcc(), "-ccbin", driver_cc, "-std=c++17", "-O2", + "--compiler-options", "-fPIC" }; + // The host compiler nvcc drives is not one mcpp resolved, so nothing + // has told it where the C library or the assembler are. Measured: with + // neither of these, NVIDIA's own crt/host_config.h stops at + // `features.h: No such file or directory`. Both are the flags mcpp + // passes to its own compiler for this target. + if (const char* sr = mcpp::toolchain_sysroot(); sr && *sr) { + front.push_back("--compiler-options"); + front.push_back(std::string("--sysroot=") + sr); + } + if (const char* bu = mcpp::toolchain_binutils_dir(); bu && *bu) { + front.push_back("--compiler-options"); + front.push_back(std::string("-B") + bu); + } + for (auto const& inc : tk->include_dirs()) front.push_back("-I" + inc); + for (auto const& a : tg.archs) { + std::string digits; + for (char c : a) if (std::isdigit(static_cast(c))) digits += c; + front.push_back("-gencode"); + front.push_back(std::format("arch=compute_{},code={}", digits, a)); + } + if (!tg.ptx.empty()) { + front.push_back("-gencode"); + front.push_back(std::format("arch=compute_{0},code=compute_{0}", tg.ptx)); + } + if (auto missing = unreachable_stage(*tk, driver_cc)) + mcpp::warning(std::format( + "nvcc cannot reach its own back-end: it invokes '{}' by name and that name " + "does not resolve on the search path it states. On the 13.x line install " + "xim:libnvvm beside xim:cuda-nvcc.", *missing).c_str()); + std::println("mcpplibs.rules.cuda: nvcc route -- {} with -ccbin {}", tk->nvcc(), driver_cc); } - std::println("example.rules.cuda: nvcc {} with -ccbin {}", nvcc, ccbin); + + // The link line gets its directories from here, not from the manifest: the + // rule resolved the payload, so the rule names where its libraries are. + for (auto const& d : tk->lib_dirs()) mcpp::link_search(d.c_str()); for (auto const& src : sources) { const auto stem = std::filesystem::path(src).stem().string(); const auto obj = opt.out_dir + "/" + stem + ".cu.o"; edge e; e.id = "cuda:" + stem; - e.description = "nvcc " + src; - e.command = { nvcc, "-c", root + "/" + src, "-o", obj, - "-ccbin", ccbin, "-std=c++17", "-O2", - "--compiler-options", "-fPIC" }; + e.description = (r == route::clang ? "clang -x cuda " : "nvcc ") + src; + e.command = front; for (auto const& inc : opt.includes) - e.command.push_back("-I" + root + "/" + inc); - for (auto const& a : opt.archs) { - // `compute_NN` is the virtual architecture the real one derives - // from; nvcc wants both halves named. - std::string digits; - for (char c : a) if (std::isdigit(static_cast(c))) digits += c; - e.command.push_back("-gencode"); - e.command.push_back(std::format("arch=compute_{},code={}", digits, a)); - } - if (!opt.ptx.empty()) { - e.command.push_back("-gencode"); - e.command.push_back(std::format("arch=compute_{0},code=compute_{0}", opt.ptx)); - } + e.command.push_back("-I" + (std::filesystem::path(inc).is_absolute() + ? inc : root + "/" + inc)); + e.command.insert(e.command.end(), { "-c", root + "/" + src, "-o", obj }); e.inputs = { root + "/" + src }; e.outputs = { obj }; out.push_back(std::move(e)); @@ -183,13 +586,10 @@ inline std::vector plan(std::span sources, options opt } inline bool submit(std::span edges) { - if (edges.empty()) return false; for (auto const& e : edges) { mcpp::action a; a.id = e.id.c_str(); - // `object`, not `source`: nvcc produces the linkable artifact itself. - // What the role names is what the output IS, not how it was made. - a.role = "object"; + a.role = "object"; // the linkable artifact itself a.description = e.description.c_str(); for (auto const& c : e.command) a.arg(c.c_str()); for (auto const& i : e.inputs) a.input(i.c_str()); @@ -199,9 +599,21 @@ inline bool submit(std::span edges) { return true; } -inline bool compile(std::span sources, options opt = {}) { +// Everything from the manifest: the architectures from `[build] accel`, the +// sources from the constrained glob in `[build] sources`. A build that asks +// for no accelerator has no device sources and nothing to do here -- that is +// the CPU-only variant, and the seam's fallback carries it. +inline bool compile(options opt = {}) { + if (!*mcpp::accel()) return true; + const auto sources = device_sources(); + if (sources.empty()) { + mcpp::warning("[build] accel names a device but no constrained glob matched a device " + "source; nothing was compiled for it"); + return true; + } auto edges = plan(sources, std::move(opt)); + if (edges.empty()) return false; return submit(edges); } -} // namespace example::rules::cuda +} // namespace mcpplibs::rules::cuda diff --git a/mcpp.toml b/mcpp.toml index 9d8e7e022..2d6ed50c8 100644 --- a/mcpp.toml +++ b/mcpp.toml @@ -1,6 +1,6 @@ [package] name = "mcpp" -version = "2026.9.5.1" +version = "2026.9.5.2" description = "Modern C++ build & package management tool" license = "Apache-2.0" authors = ["mcpp-community"] diff --git a/modules/buildmcpp/src/directives.cppm b/modules/buildmcpp/src/directives.cppm index 30cf3d736..69df97350 100644 --- a/modules/buildmcpp/src/directives.cppm +++ b/modules/buildmcpp/src/directives.cppm @@ -122,6 +122,19 @@ enum class Slot : std::size_t { // A sentence for the USER. Not a build input at all — see Scope::Advisory // for why this could not be folded into any existing slot. Warnings, + // ⭐ A CLAIM ABOUT THE MACHINE, OR ABOUT WHAT THIS PACKAGE NEEDS OF IT. + // + // `fact` carries `=`: something the program established + // about the machine, by whatever means the package owns (a driver's + // version, read from the driver's own library). `floor` carries + // ` >= `: what the package needs of that quantity. The + // engine compares the two before anything is compiled and refuses with + // both values when the floor is unmet (`version-floor-unmet`); a floor + // nobody stated a fact for is silent. Neither string means anything to + // this file -- the name is data flowing through -- which is what keeps + // vendor knowledge in the package that has it and out of the engine. + Facts, + Floors, Count }; inline constexpr std::size_t kSlotCount = static_cast(Slot::Count); @@ -162,6 +175,13 @@ enum class Scope { // it is still wrong: a re-run key feeds a MACHINE decision, an advisory // feeds a person. Advisory, + // A statement the engine COMPARES at prepare time. It reaches no compile + // line, no link line, no source set and no person directly; the verdict + // of the comparison does. Persisted, so a cached run replays the claim -- + // right for a floor, and the reason a program stating a FACT about the + // machine must also declare what would change it (`rerun_if_changed` on + // the file the fact was read from), or the fact outlives the machine. + Claim, }; // How the raw wire value is normalized before it is stored. Applied ONCE, at @@ -197,7 +217,7 @@ struct Def { int sinceProtocol; }; -inline constexpr std::array kTable{{ +inline constexpr std::array kTable{{ // wire tag slot scope transform must missingPrefix missingSuffix since {"cxxflag", "cxxflag", Slot::CxxFlags, Scope::PackagePrivate, Transform::Verbatim, false, "", "", 1}, {"cflag", "cflag", Slot::CFlags, Scope::PackagePrivate, Transform::Verbatim, false, "", "", 1}, @@ -283,6 +303,10 @@ inline constexpr std::array kTable{{ // the unknown-tag path and discards the whole record. {"warning", "warning", Slot::Warnings, Scope::Advisory, Transform::Verbatim, false, "", "", 5}, {"action", "action", Slot::Actions, Scope::GraphNode, Transform::Verbatim, false, "", "", 1}, + // The probe channel: a rule package measures, the engine compares. See + // Slot::Facts for the shape of each value. + {"fact", "fact", Slot::Facts, Scope::Claim, Transform::Verbatim, false, "", "", 7}, + {"floor", "floor", Slot::Floors, Scope::Claim, Transform::Verbatim, false, "", "", 7}, }}; // ── Collected output of one run ──────────────────────────────────────────── @@ -747,6 +771,19 @@ void apply(mcpp::manifest::Manifest& m, const Directives& d) { for (auto const& p : d.at(Slot::IncludeDirsAfter)) bc.includeDirsAfter.emplace_back(p); + // Claims join the runtime declarations the manifest could have carried + // itself, so the version-floor check in prepare reads ONE list and never + // learns which spelling a claim arrived in. + for (auto const& f : d.at(Slot::Facts)) + m.runtimeConfig.provides.push_back(f); + for (auto const& fl : d.at(Slot::Floors)) { + mcpp::manifest::RuntimeRequirement req; + req.kind = "version-floor"; + req.value = fl; + req.phase = "build"; + m.runtimeConfig.requirements.push_back(std::move(req)); + } + // Build-graph nodes. Decoded here rather than at parse time so the cache // stores the payload verbatim and a replay is byte-identical to a run. for (auto const& payload : d.at(Slot::Actions)) { diff --git a/modules/buildmcpp/src/program_protocol.cppm b/modules/buildmcpp/src/program_protocol.cppm index 058eaaa92..8cbc98aca 100644 --- a/modules/buildmcpp/src/program_protocol.cppm +++ b/modules/buildmcpp/src/program_protocol.cppm @@ -56,7 +56,11 @@ export namespace mcpp::build::program_protocol { // cost as v5's: a package calling `mcpp::flash()` fails on an older engine at // the build.mcpp COMPILE, because that engine's bundled module has no such // function, not through a protocol refusal. -inline constexpr int kProtocolVersion = 6; +// v7: adds `fact` and `floor` -- the probe channel. A rule package states a +// fact about the machine and the floor it needs of it; the engine compares +// them before compiling. Same cost as v5's: a package calling `mcpp::fact()` +// fails on an older engine at the build.mcpp COMPILE, not through a refusal. +inline constexpr int kProtocolVersion = 7; // ── Cache-format epoch ───────────────────────────────────────────────────── // diff --git a/modules/manifest/src/toml.cppm b/modules/manifest/src/toml.cppm index fc4279f34..0333df428 100644 --- a/modules/manifest/src/toml.cppm +++ b/modules/manifest/src/toml.cppm @@ -485,6 +485,14 @@ std::expected parse_string(std::string_view content, // find_disallowed_array_of_tables above. static constexpr std::string_view kAllowedArraysOfTables[] = { "build.flags", + // A `sources` list whose every entry is a table (`[{ glob, accel }]`) + // has the same Value shape as `[[build.sources]]`; the reader + // type-checks every entry, so nothing is silently taken. The + // conditional axis is allowlisted for the same reason: its reader + // refuses a table entry with a message that says where it belongs, + // which this guard's generic sentence would pre-empt. + "build.sources", + "target.*.build.sources", "features.*.flags", // #253 — the middle segment is the feature name "target.*.build.flags", // #258 — middle segment is the cfg predicate "runtime.requirements", @@ -594,8 +602,48 @@ std::expected parse_string(std::string_view content, // mean "compile nothing", and only the key's existence can say that (see // BuildConfig::sourcesDeclared). Set from either spelling, because the // legacy one has to be able to express it too. - if (auto v = doc->get_string_array("build.sources")) { - m.buildConfig.sources = *v; + // A mixed list: plain globs, and tables that carry a constraint. Read + // element by element rather than through get_string_array, which answers + // "not a string array" for a list with one table in it -- and then the + // key would count as absent, the default glob would apply, and the + // constrained entry would vanish without a word. + if (auto* sv = doc->get("build.sources")) { + if (!sv->is_array()) { + return std::unexpected(error(origin, + "[build].sources must be an array of globs and/or inline tables " + "({ glob = \"...\", accel = \"...\" })")); + } + for (auto& ev : sv->as_array()) { + if (ev.is_string()) { m.buildConfig.sources.push_back(ev.as_string()); continue; } + if (!ev.is_table()) { + return std::unexpected(error(origin, + "[build].sources entries must be strings or inline tables " + "with a `glob` key")); + } + mcpp::manifest::BuildConfig::SourceConstraint sc; + for (auto& [k, v] : ev.as_table()) { + bool ok = false; + if (k == "glob") { ok = v.is_string(); if (ok) sc.glob = v.as_string(); } + else if (k == "accel") { ok = v.is_string(); if (ok) sc.accel = v.as_string(); } + if (!ok) { + return std::unexpected(error(origin, std::format( + "[build].sources: invalid key '{}' in a table entry " + "(expected glob = \"...\" and optionally accel = \"...\")", k))); + } + } + if (sc.glob.empty()) { + return std::unexpected(error(origin, + "[build].sources: a table entry is missing its `glob` key")); + } + if (sc.glob.starts_with("!")) { + return std::unexpected(error(origin, std::format( + "[build].sources: a constrained entry cannot be an exclusion " + "('{}'); write the exclusion as a plain string", sc.glob))); + } + m.buildConfig.sources.push_back(sc.glob); + if (!sc.accel.empty()) + m.buildConfig.sourceConstraints.push_back(std::move(sc)); + } m.buildConfig.sourcesDeclared = true; } if (auto v = doc->get_string_array("modules.sources")) { @@ -2451,6 +2499,17 @@ std::expected parse_string(std::string_view content, read_list("cflags", cc.inputs.cflags); read_list("cxxflags", cc.inputs.cxxflags); read_list("ldflags", cc.inputs.ldflags); + // A constrained entry (`{ glob, accel }`) is refused here rather + // than skipped: read_list keeps strings only, and a table that + // silently vanished would be a device glob nobody ever narrowed. + if (auto f = bt.find("sources"); f != bt.end() && f->second.is_array()) + for (auto& v : f->second.as_array()) + if (!v.is_string()) { + return std::unexpected(error(origin, + "[target.'cfg(...)'.build].sources entries must be " + "plain globs; a constrained entry ({ glob, accel }) " + "belongs in [build].sources")); + } read_list("sources", cc.inputs.sources); // #296: package-level macros are a build input like any other, // so the cfg axis carries them too — a platform-only macro diff --git a/modules/manifest/src/types.cppm b/modules/manifest/src/types.cppm index 519f832d2..7cda33669 100644 --- a/modules/manifest/src/types.cppm +++ b/modules/manifest/src/types.cppm @@ -550,6 +550,26 @@ struct BuildConfig : BuildInputs { // the "main" feature) without it being linked by default — see // .agents/docs/2026-06-25-gtest-main-feature-and-add-dev-design.md. std::map> featureSources; + // `[build] sources` entries written as a table: + // + // sources = ["src/**/*.cppm", { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }] + // + // The glob ALSO appears in `sources`, so every reader that walks the plain + // list sees it; this carries the constraint that decides whether it + // applies to a given build. Resolved in prepare_build: the glob must match + // at least one file (an empty match is a typo, not a no-op); when the + // build asks for no accelerator the glob is excluded, which is how + // `--no-accel` yields the CPU-only variant of a project; and when it does + // ask for one, the constraint must lie within what the build targets, or + // the build is refused naming both. Device-kind files the effective source + // set matches are handed to the package's build program + // (MCPP_DEVICE_SOURCES) rather than compiled by the engine, which has no + // rule for them: that is the rule package's business. + struct SourceConstraint { + std::string glob; + std::string accel; // wire form, the mcpp.pack.abi_tag grammar + }; + std::vector sourceConstraints; // feature name → package-owned preprocessor defines (e.g. "-DEIGEN_USE_BLAS"). // Feature System v2 Stage 1: when the feature is active these are appended to // the package's compile flags alongside the automatic -DMCPP_FEATURE_ diff --git a/modules/versioning/src/version.cppm b/modules/versioning/src/version.cppm index 07110adb2..c035f79ee 100644 --- a/modules/versioning/src/version.cppm +++ b/modules/versioning/src/version.cppm @@ -31,6 +31,6 @@ import std; export namespace mcpp { -inline constexpr std::string_view MCPP_VERSION = "2026.9.5.1"; +inline constexpr std::string_view MCPP_VERSION = "2026.9.5.2"; } // namespace mcpp diff --git a/src/build/build_program.cppm b/src/build/build_program.cppm index 3dd29a335..f4274f00c 100644 --- a/src/build/build_program.cppm +++ b/src/build/build_program.cppm @@ -43,6 +43,29 @@ struct BuildProgramEnv { // hostprogram::toolchain_dir / sysroot_dir for why declaring was wrong. std::string toolchainDir; std::string targetSysroot; + // ⭐⭐ THE TWO ANSWERS A SECOND COMPILER NEEDS AND CANNOT DERIVE. + // + // `toolchainSysroot` is the `--sysroot` mcpp passes to its own compiler and + // `toolchainBinutilsDir` the directory it names with `-B`; either is empty + // when mcpp passes none. They are not the same question as + // `targetSysroot`, which is a TIER fact (a bare-metal target's own C + // library payload, empty on a hosted target) — these two are ENVIRONMENT + // facts, and on a hosted subos both are non-empty precisely because the C + // library is not at `/usr/include` and the assembler is not at `/usr/bin`. + // + // ⚠️ Measured 2026-09-05 on the CUDA example. `nvcc` refuses a libc++ host + // compiler and fails on GCC 16's ``, so its rule package + // resolves a second host compiler from a declared payload. That compiler + // is not one mcpp resolved, so nothing tells it where anything is, and the + // first `#include` in NVIDIA's own `crt/host_config.h` fails: + // + // host_config.h:218: fatal error: features.h: No such file or directory + // + // Every rule package driving a compiler mcpp did not resolve has the same + // gap — `hipcc`, `-fsycl-host-compiler`, a generator that compiles what it + // emits — so the answer belongs to the engine and is stated once here. + std::string toolchainSysroot; + std::string toolchainBinutilsDir; // ⭐⭐ WHICH COMPILER RESOLVED — "gcc" | "clang" | "msvc" | "". // // A package should never have to guess this, and until this field existed @@ -77,6 +100,20 @@ struct BuildProgramEnv { std::string targetLibc; // "picolibc-riscv" | "" (zero-libc tier) std::string profile; // effective profile name (dev/release/…) std::vector features; // active feature closure of the package + // The device axis of this build, in the wire form `mcpp.pack.abi_tag` + // reads (`cuda12.9+{sm_89} ptx>=89`); empty when the build asks for no + // accelerator. Already resolved -- `--accel` / `--no-accel` over + // `[build] accel` -- so a rule package derives its own spelling + // (`-gencode`, `--offload-arch`) from here and the architecture set is + // written once, in the manifest, and never again in a build program. + std::string accel; + // The device-kind sources (`.cu`, `.hip`, ...) this package's effective + // source set matches, package-root-relative with `/` separators, one per + // line. The engine has no compile rule for them and hands the list to the + // build program, where the rule package the package imports turns each + // one into an `mcpp::action`. Already narrowed: a glob whose `accel` + // constraint the build does not satisfy contributes nothing. + std::vector deviceSources; // Artifact home (bin/cache/out). Empty → /target/.build-mcpp (the // root-project default). Dependencies MUST point this into the CONSUMING // project's tree — a registry package root is shared and may be read-only. @@ -420,12 +457,23 @@ contract_env(const fs::path& root, const fs::path& outDir, const BuildProgramEnv // absent variable would make the answer depend on whatever the parent // process happened to export. e.emplace_back("MCPP_TOOLCHAIN_DIR", env.toolchainDir); + e.emplace_back("MCPP_TOOLCHAIN_SYSROOT", env.toolchainSysroot); + e.emplace_back("MCPP_TOOLCHAIN_BINUTILS_DIR", env.toolchainBinutilsDir); e.emplace_back("MCPP_COMPILER", env.compilerId); e.emplace_back("MCPP_TARGET_SYSROOT", env.targetSysroot); e.emplace_back("MCPP_TARGET_BUILTINS_LIB", env.targetBuiltinsLib); e.emplace_back("MCPP_TARGET_LIBC_PROFILE", env.targetLibcProfile); e.emplace_back("MCPP_TARGET_LIBC", env.targetLibc); e.emplace_back("MCPP_PROFILE", env.profile); + e.emplace_back("MCPP_ACCEL", env.accel); + { + std::string joined; + for (auto const& d : env.deviceSources) { + if (!joined.empty()) joined += '\n'; + joined += d; + } + e.emplace_back("MCPP_DEVICE_SOURCES", joined); + } e.emplace_back("MCPP_OUT_DIR", outDir.string()); e.emplace_back("MCPP_MANIFEST_DIR", root.string()); std::string csv; @@ -699,7 +747,7 @@ std::expected run_build_program( compilerIdentity += "\nbuild-program-link="; compilerIdentity += muslStaticHelper ? "musl-static-v1" : mingwStaticHelper ? "mingw-static-v1" - : "default-v1"; + : "default-v2"; // v2: DT_RPATH on Linux std::string programHash = mcpp::toolchain::hash_file(src); std::string compilerHash = mcpp::toolchain::hash_string(compilerIdentity); @@ -1000,6 +1048,23 @@ std::expected run_build_program( // compile/precompile commands, where a link flag has no business (and for // Clang would perturb the default PIC/PIE codegen of mcpp.o). if (staticHostHelper) compileArgv.push_back(std::string(dial.staticRuntime)); + // ⚠️ A DYNAMIC HELPER ON LINUX GETS `DT_RPATH`, NOT `DT_RUNPATH`. + // + // The driver's default is the new tag, and a runpath is consulted only for + // the helper's OWN needed libraries. A build program that opens a host + // library at run time -- a rule package reading a driver's version through + // the driver itself -- then fails one hop later, because that library's + // own dependencies (`libdl.so.2`, `libpthread.so.0`) are looked up without + // the helper's search path and the payload loader has no default that + // reaches them. Measured: `dlopen("/lib/libcuda.so.1")` from a + // build.mcpp answered `libdl.so.2: cannot open shared object file` while + // the very same directories sat in the helper's RUNPATH. The artifacts + // mcpp links carry DT_RPATH for this reason (loader_contract's Rpath tag); + // the helper now does too. Driver-only spelling: the helper is always + // linked through the compiler driver, never through the linker directly. + if (!staticHostHelper && !msvcHost + && !mcpp::platform::is_windows && !mcpp::platform::is_macos) + compileArgv.push_back("-Wl,--disable-new-dtags"); if (msvcHost) { // /Fe: takes its value attached, not as a separate argv token. compileArgv.push_back(std::string(dial.outputExePrefix) + bin.string()); diff --git a/src/build/execute.cppm b/src/build/execute.cppm index d78e1386c..eb3fea845 100644 --- a/src/build/execute.cppm +++ b/src/build/execute.cppm @@ -1154,13 +1154,15 @@ export int list_runners(const std::string& package_filter, // for a probe. Reporting them without the axis that // selects them would answer a question nobody asked. const std::string& features = {}, - const std::string& profile = {}) { + const std::string& profile = {}, + const std::string& accel = {}) { mcpp::build::BuildOverrides ov; ov.package_filter = package_filter; ov.cache_mode = no_cache ? std::string("off") : cache_mode; ov.target_triple = target_triple; ov.features = features; ov.profile = profile; + ov.accel = accel; // Reporting what `mcpp run` would do means resolving what `mcpp run` // resolves, tool tiers included — otherwise this command would list a // runner whose program it had declined to install. @@ -1540,7 +1542,10 @@ export int build_run_target(const std::optional& targetName, // the board arrives. Without this it was the one // scenario the design's own example could not run. const std::string& features = {}, - const std::string& profile = {}) { + const std::string& profile = {}, + // The device axis. `--no-accel` arrives as the + // "(none)" sentinel, as it does for `build`. + const std::string& accel = {}) { // mcpp#225 (E2): reuse the resolved build cache when it's still fresh, // skipping prepare_build's toolchain resolution + modgraph scan // entirely — mirrors cmd_build's try_fast_build fast path. The cached @@ -1561,7 +1566,7 @@ export int build_run_target(const std::optional& targetName, // whichever feature set and profile the last build used; taking it // here would silently ignore the flag, which is the same reason // `--cache` and `--profile` bypass it in `cmd_build`. - && features.empty() && profile.empty() + && features.empty() && profile.empty() && accel.empty() // ⚠️⚠️ THE FAST PATH IS `run`'s, AND ONLY `run`'s. // // It exec's the cached artefact directly — that IS its definition — so @@ -1588,6 +1593,7 @@ export int build_run_target(const std::optional& targetName, ov.target_triple = target_triple; ov.features = features; ov.profile = profile; + ov.accel = accel; // This verb executes what it builds, so the `when = "run"` tool tier is // part of what has to exist. `mcpp build` does not set it, which is the // whole of the difference the tier buys. diff --git a/src/build/hostprogram.cppm b/src/build/hostprogram.cppm index 511266b18..d3d4cb4a9 100644 --- a/src/build/hostprogram.cppm +++ b/src/build/hostprogram.cppm @@ -102,6 +102,28 @@ inline void run_exclusive() { std::printf("mcpp:run-exclusive // "resolved". mcpp replays it on every hit. inline void warning(const char* message) { std::printf("mcpp:warning=%s\n", message); } +// ── The probe channel (mcpp 2026.9.5.2+) ──────────────────────────────────── +// +// A rule package is the thing that knows how to ask a machine what it has -- +// which library to open, which function to call -- and the engine is the +// thing that must not. So the package MEASURES and the engine COMPARES: +// +// mcpp::fact("cuda.driver", "12.4"); // what this machine has +// mcpp::floor("cuda.driver >= 12.0"); // what this package needs +// +// Before anything is compiled the engine refuses when a floor is unmet and +// names both values (`mcpp why toolchain --format json` reports the reason +// `version-floor-unmet`). A floor nobody stated a fact for is silent: not +// knowing is not failing. +// +// ⚠️ A fact is persisted with the program's other output and replayed on a +// cache hit. State what would change it -- `rerun_if_changed` on the library +// the version was read from -- or the fact outlives the machine it described. +inline void fact(const char* name, const char* version) { + std::printf("mcpp:fact=%s=%s\n", name, version); +} +inline void floor(const char* spec) { std::printf("mcpp:floor=%s\n", spec); } + // The memory layout for a freestanding link. Reaches the CONSUMER's link line // (like link_lib/link_search, unlike include_dir), because the package that // knows a board's layout is not the package being built. @@ -214,6 +236,20 @@ inline const char* target_arch() { return env_or("MCPP_TARGET_A inline const char* target_env() { return env_or("MCPP_TARGET_ENV"); } inline const char* host() { return env_or("MCPP_HOST"); } inline const char* profile() { return env_or("MCPP_PROFILE"); } +// The device axis of this build: `cuda12.9+{sm_89} ptx>=89`, or "" when the +// build asks for no accelerator. Already resolved (`--accel` / `--no-accel` +// over `[build] accel`), so a rule package derives its architecture flags +// from HERE and the set is written once, in the manifest. What the string +// means beyond "backend, version, architectures, floor" is the package's +// business: the engine never learns what `sm_89` is. +inline const char* accel() { return env_or("MCPP_ACCEL"); } +// The device-kind sources (`.cu`, `.hip`, ...) this package's `sources` match +// under the current accel, package-root-relative, one per line, "" when there +// are none. The engine compiles none of them; the rule package this program +// imports turns each into an `mcpp::action`. Already narrowed: a glob written +// as `{ glob = "...", accel = "..." }` whose constraint the build does not +// satisfy contributes nothing, so `--no-accel` yields an empty list. +inline const char* device_sources() { return env_or("MCPP_DEVICE_SOURCES"); } inline const char* out_dir() { return env_or("MCPP_OUT_DIR"); } // Where the TOOLCHAIN mcpp resolved for this build lives — the payload root, @@ -228,6 +264,21 @@ inline const char* out_dir() { return env_or("MCPP_OUT_DIR" // `[toolchain]` actually resolved. inline const char* toolchain_dir() { return env_or("MCPP_TOOLCHAIN_DIR"); } +// The two flags mcpp passes to its own compiler: the `--sysroot` and the +// directory it names with `-B`. Either is empty when mcpp passes none. +// +// ⭐ These are for A SECOND COMPILER — one this rule package runs and mcpp did +// not resolve. Such a compiler starts with no idea where anything is, and on a +// subos the C library is not at `/usr/include` and the assembler is not at +// `/usr/bin`; the first `#include` it reaches then fails on `features.h`. +// Forwarding these two makes it see what mcpp's own compiler sees. +// +// ⚠️ NOT `sysroot_dir()`, four lines down. That one answers a question about +// the TARGET's tier and is empty on a hosted target, which is exactly the case +// this pair exists for. +inline const char* toolchain_sysroot() { return env_or("MCPP_TOOLCHAIN_SYSROOT"); } +inline const char* toolchain_binutils_dir() { return env_or("MCPP_TOOLCHAIN_BINUTILS_DIR"); } + // Which compiler resolved: "gcc", "clang", "msvc", or "" if none did. // // ⭐ Ask this rather than inferring it from `toolchain_dir()`. The two questions diff --git a/src/build/prepare.cppm b/src/build/prepare.cppm index 2416aa452..e33d2871e 100644 --- a/src/build/prepare.cppm +++ b/src/build/prepare.cppm @@ -19,6 +19,7 @@ import std; import mcpp.targetside; import mcpp.diag; import mcpp.build.refusal; +import mcpp.build.version_floor; import mcpp.home; import mcpp.platform.axis; import mcpp.libs.json; @@ -36,6 +37,8 @@ import mcpp.toolchain.dialect; import mcpp.toolchain.fingerprint; import mcpp.toolchain.msvc; import mcpp.toolchain.registry; +import mcpp.toolchain.linkmodel; +import mcpp.toolchain.gcc; // For `resolve_version_match` / `list_installed_versions`: a bare compiler // family named by the dependency graph resolves to a concrete version through // exactly the path `mcpp toolchain default ` uses. @@ -1017,6 +1020,16 @@ void fill_target_build_env(mcpp::build::BuildProgramEnv& e, e.targetLibc = tc ? tc->targetSysrootPkg : std::string{}; if (!tc) return; + // The two flags mcpp passes to ITS OWN compiler, so a rule package driving + // a second compiler passes the same two. Both read from the single + // producer that already decides them for the engine's own command lines — + // `resolve_link_model` for the sysroot, `gcc::binutils_prefix_dir` for the + // `-B` — rather than a fifth re-derivation of either. + if (auto lm = mcpp::toolchain::resolve_link_model(*tc); + lm.mode == mcpp::toolchain::CLibMode::Sysroot) + e.toolchainSysroot = lm.sysroot.string(); + e.toolchainBinutilsDir = mcpp::toolchain::gcc::binutils_prefix_dir(*tc).string(); + // The C LIBRARY's sub-directory for this ISA profile, from the freestanding // table — the same single read point the compile flags use. // @@ -2389,8 +2402,43 @@ prepare_build(bool print_fingerprint, // canonicalized. Reading it before that point would silently fall back to // the host for any project that sets its target in the manifest rather // than on the command line. - const auto targetPlatform = mcpp::platform::TargetPlatform::for_os( - cfgpred::context_for(overrides.target_triple).os); + // ── The device axis, resolved ONCE ──────────────────────────────────── + // + // `--accel` / `--no-accel` over `[build] accel`. `--no-accel` arrives as the + // sentinel "(none)", which parse_accel reads as nothing, and printing the + // parsed form back normalises the spelling -- so every reader below sees + // one string, and a build program sees the same one in MCPP_ACCEL. Read + // at call time rather than captured: a `[target.'cfg(...)'.build]` section + // may set `accel`, and the merge that applies it runs a few lines down. + // + // ⚠️⚠️ "NO ACCELERATOR" IS THE EMPTY STRING HERE, NOT `accel_str`'s "(none)". + // + // `accel_str` is a DISPLAY function: it prints `(none)` for an empty set so + // an ABI tag reads as a sentence. Handing that spelling on as a value made + // two readers wrong at once. A build program saw `MCPP_ACCEL=(none)` while + // the manual promised an empty string, so a rule package asking "is there + // an accelerator" got a yes and a backend named `(none)`; and the + // fingerprint's own guard, `if (!accel.empty())`, was true for every + // project on earth, appending `#accel=(none)` to builds that had asked for + // nothing. Measured 2026-09-05 with a build program that wrote the value to + // a file, which is the only way to see it -- a program's stdout is shown + // only when it fails. + auto resolvedAccel = [&]() -> std::string { + const auto sets = mcpp::pack::parse_accel( + overrides.accel.empty() ? m->buildConfig.accel : overrides.accel); + return sets.empty() ? std::string{} : mcpp::pack::accel_str(sets); + }; + // The cfg context, with the accelerator layer filled from the resolved + // accel's backend names. `cfg(accelerator = "cuda")` is a membership test + // over these (prepare_inputs::Ctx::layer_matches); before this the field + // was declared, documented, and never written, so the key matched nothing. + auto cfgCtx = [&]() { + auto c = cfgpred::context_for(overrides.target_triple); + for (auto const& set : mcpp::pack::parse_accel(resolvedAccel())) + c.accelerators.push_back(set.backend); + return c; + }; + const auto targetPlatform = mcpp::platform::TargetPlatform::for_os(cfgCtx().os); // ── L1: merge conditional [target.'cfg(...)'] sections ─────────────────── // Evaluated now (target resolved) against the resolved target — the @@ -2407,7 +2455,7 @@ prepare_build(bool print_fingerprint, // package's half of the one funnel, not a special case: every package is // merged exactly once, immediately before it is captured into `packages[]`. if (!m->conditionalConfigs.empty()) { - merge_conditional_config(*m, cfgpred::context_for(overrides.target_triple)); + merge_conditional_config(*m, cfgCtx()); } // `[build].defines` must reach the scanner (P1689) and the compile edge, // and must participate in the fingerprint. Fold before dependency @@ -4529,7 +4577,7 @@ prepare_build(bool print_fingerprint, // pass through loadVersionDep. if (!manifest->conditionalConfigs.empty()) { merge_conditional_config(*manifest, - cfgpred::context_for(overrides.target_triple)); + cfgCtx()); } fold_build_defines_into_flags(manifest->buildConfig); @@ -5936,7 +5984,7 @@ prepare_build(bool print_fingerprint, // snapshot this manifest's flags/sources into `packages[]`. if (!dep_manifest->conditionalConfigs.empty()) { merge_conditional_config(*dep_manifest, - cfgpred::context_for(overrides.target_triple)); + cfgCtx()); } fold_build_defines_into_flags(dep_manifest->buildConfig); } else { @@ -6477,6 +6525,65 @@ prepare_build(bool print_fingerprint, // the question it answers is different: capProviders asks "can this // requirement be satisfied", this asks "can these two coexist at all". std::map> capExclusive; + // Callable twice: once here, for what the manifests and the + // dependencies' build programs declared, and once more after the + // root's build program has run -- a rule package it imports states + // its facts and floors from there (`mcpp::fact` / `mcpp::floor`), + // and a check that ran only before it would never see them. + // package name -> device-kind sources of its effective source set, filled + // by the narrowing pass after feature application and read at both + // build-program run sites (MCPP_DEVICE_SOURCES). + // Keyed by the package's ROOT DIRECTORY, not by its name. Two packages in + // one graph may share a bare name and differ only by namespace — that is + // what namespaces are for — and a name key would hand one package's + // device sources to the other's build program with nothing reporting it. + std::map> deviceSourcesByPackage; + auto checkVersionFloors = [&]() -> std::optional { + std::map> facts; // name -> (version, who) + for (std::size_t pi = 0; pi < packages.size(); ++pi) { + // The root's claims live in *m: its build program mutates + // *m, and packages[0] is a snapshot taken before it ran. + const auto& mf = pi == 0 ? *m : packages[pi].manifest; + const auto who = mf.package.name; + for (auto const& entry : mf.runtimeConfig.provides) { + auto fact = mcpp::build::parse_version_fact(entry); + if (fact.valid()) facts.emplace(fact.name, std::pair{fact.version, who}); + } + } + for (std::size_t pi = 0; pi < packages.size(); ++pi) { + // The root's claims live in *m: its build program mutates + // *m, and packages[0] is a snapshot taken before it ran. + const auto& mf = pi == 0 ? *m : packages[pi].manifest; + const auto who = mf.package.name; + for (auto const& req : mf.runtimeConfig.requirements) { + if (req.kind != "version-floor") continue; + auto floor = mcpp::build::parse_version_floor(req.value); + if (!floor.valid()) { + return std::format( + "`{}` declares a version-floor requirement mcpp " + "cannot read: '{}'.\n" + " The shape is ` >= `, e.g. " + "`cuda.driver >= 12.0`.", who, req.value); + } + auto it = facts.find(floor.name); + if (it == facts.end()) continue; // nobody stated it + auto met = mcpp::build::version_at_least(it->second.first, + floor.version); + if (!met || *met) continue; + refusal::record(refusal::Code::VersionFloorUnmet); + return std::format( + "`{}` requires {} >= {}, and this machine has {}.\n" + " stated by: {}\n" + " This is checked before anything is compiled " + "because the failure it prevents is not:\n" + " a build against too-new a runtime links " + "cleanly and fails at first use.", + who, floor.name, floor.version, it->second.first, + it->second.second); + } + } + return std::nullopt; + }; { auto sanitize = [](std::string f) { for (auto& c : f) @@ -6803,6 +6910,83 @@ prepare_build(bool print_fingerprint, activeFeaturesByPackage[i] = feature_closure(packages[i].manifest, req, depDefaultFeatures); } + + // ── Constrained source globs: narrow to what this build targets ──── + // + // A `{ glob = "...", accel = "..." }` entry in `[build] sources` says + // what its files are FOR. Three outcomes, all decided here and none in + // the scanner, which keeps reading a plain list of globs: + // + // - the glob matches nothing: refused, naming the glob. An empty + // match is a typo or a moved directory, not a no-op, and the + // failure it would otherwise become is a kernel that is never + // compiled and a link that resolves nothing. + // - the build asks for no accelerator: the glob is EXCLUDED, with the + // same `!` mechanism feature gates use -- removing the string is not + // enough when a broader glob (the default `src/**`) covers the same + // files. This is how `--no-accel` yields the CPU-only variant. + // - the build asks for one: the constraint must lie within it, or the + // build is refused naming both. A file compiled for sm_89 under a + // build that targets sm_80 is not a variant, it is a mismatch. + // + // Device-kind files the effective set still matches are collected per + // package for the build program (MCPP_DEVICE_SOURCES); the engine has + // no compile rule for them and never will. + { + const auto buildAccel = mcpp::pack::parse_accel(resolvedAccel()); + for (std::size_t i = 0; i < packages.size(); ++i) { + auto& pkg = packages[i]; + auto& bc = pkg.manifest.buildConfig; + std::set excludedGlobs; + for (auto const& sc : bc.sourceConstraints) { + const auto hits = mcpp::modgraph::expand_glob(pkg.root, sc.glob); + if (hits.empty()) { + return std::unexpected(std::format( + "`{}`: [build] sources entry '{}' (accel = \"{}\") matches no file.\n" + " A constrained glob names the files a device build needs; an\n" + " empty match would leave nothing to compile for that device\n" + " and say so only at the link, or never.", + pkg.manifest.package.name, sc.glob, sc.accel)); + } + if (buildAccel.empty()) { excludedGlobs.insert(sc.glob); continue; } + const auto want = mcpp::pack::parse_accel(sc.accel); + if (!mcpp::pack::accel_accepts(buildAccel, want)) { + refusal::record(refusal::Code::AccelMismatch); + return std::unexpected(std::format( + "`{}`: [build] sources entry '{}' is constrained to accel \"{}\",\n" + " which this build does not cover.\n" + " this build targets: {}\n" + " fix: build with `--accel` covering it, or `--no-accel` to\n" + " leave every constrained glob out (the CPU-only variant).", + pkg.manifest.package.name, sc.glob, + mcpp::pack::accel_str(want), + mcpp::pack::accel_str(buildAccel))); + } + } + for (auto const& g : excludedGlobs) { + bc.sources.push_back("!" + g); + pkg.manifest.modules.sources.push_back("!" + g); + } + // The device-kind files the EFFECTIVE set matches, for the + // build program. Exclusions are honoured the way the scanner + // honours them: positives first, then `!` entries removed. + const auto extTable = mcpp::extension_table_for(bc.moduleExtensions); + std::set matched, dropped; + for (auto const& g : pkg.manifest.modules.sources) { + if (g.empty()) continue; + if (g[0] == '!') { for (auto& f : mcpp::modgraph::expand_glob(pkg.root, g.substr(1))) dropped.insert(f); } + else if (!std::filesystem::path(g).is_absolute()) + for (auto& f : mcpp::modgraph::expand_glob(pkg.root, g)) matched.insert(f); + } + std::vector device; + for (auto const& f : matched) { + if (dropped.contains(f)) continue; + if (mcpp::classify(f, extTable) != mcpp::SourceKind::Device) continue; + device.push_back(f.lexically_relative(pkg.root).generic_string()); + } + deviceSourcesByPackage[pkg.root.string()] = std::move(device); + } + } activeFeaturesByPackage.resize(packages.size()); // ── #355: HOST tool provisioning ──────────────────────────────────── @@ -7346,12 +7530,16 @@ prepare_build(bool print_fingerprint, }; mcpp::build::BuildProgramEnv bpEnv; bpEnv.targetTriple = resolvedTargetCanonical; - // The payload ROOT (not the driver), the target's C library, and the - // three answers that keep a board package from hardcoding a toolchain - // or a libc. All four in one call — see fill_target_build_env. - fill_target_build_env(bpEnv, tc ? &*tc : nullptr); - bpEnv.toolsBin = projectSubosBin; + // The payload ROOT (not the driver), the target's C library, and + // the three answers that keep a board package from hardcoding a + // toolchain or a libc. All four in one call — see + // fill_target_build_env. + fill_target_build_env(bpEnv, tc ? &*tc : nullptr); + bpEnv.toolsBin = projectSubosBin; bpEnv.profile = effectiveProfile; + bpEnv.accel = resolvedAccel(); + if (auto dit = deviceSourcesByPackage.find(pkg.root.string()); dit != deviceSourcesByPackage.end()) + bpEnv.deviceSources = dit->second; bpEnv.features = feature_closure(pkg.manifest, req, depDefaultFeatures); bpEnv.artifactsDir = workRoot / "target" / ".build-mcpp" / "deps" / (dirSafe(pkg.manifest.package.name) + "@" + pkg.manifest.package.version); @@ -7524,6 +7712,19 @@ prepare_build(bool print_fingerprint, cap, claimers.size() == 1 ? "it" : "they", list, claimed)); } + // VERSION FLOORS. A package states what it needs of the machine; a + // package that established a fact about the machine states it. Neither + // string means anything to this code -- `cuda.driver` is data flowing + // through -- which is why a second backend needs no change here and why + // `test_runtime_contract`'s gate stays satisfied. + // + // ⚠️ A FLOOR WITH NO FACT IS SILENT. A machine that never declared what + // it has is not a machine that fails the floor; it is one nobody asked. + // Reporting a refusal there would turn "we do not know" into "no", and + // the whole reason this exists is that a wrong answer is worse than no + // answer. + if (auto err = checkVersionFloors(); err) return std::unexpected(*err); + std::set boundCaps; for (auto& [cap, requirer] : capRequires) { if (!boundCaps.insert(cap).second) continue; // one diagnosis per cap @@ -8175,7 +8376,7 @@ prepare_build(bool print_fingerprint, // `pkg.manifest.buildConfig` produced a build in which every layer // predicate matched and no flag reached the compiler. if (targetSideResolved) { - auto layerCtx = cfgpred::context_for(overrides.target_triple); + auto layerCtx = cfgCtx(); layerCtx.layersKnown = true; layerCtx.compiler = resolvedTargetSide.compiler.interfaceName; layerCtx.compilerRuntime = resolvedTargetSide.compilerRuntime.interfaceName; @@ -8251,6 +8452,9 @@ prepare_build(bool print_fingerprint, fill_target_build_env(bpEnv, tc ? &*tc : nullptr); bpEnv.toolsBin = projectSubosBin; bpEnv.profile = effectiveProfile; + bpEnv.accel = resolvedAccel(); + if (auto dit = deviceSourcesByPackage.find(root->string()); dit != deviceSourcesByPackage.end()) + bpEnv.deviceSources = dit->second; // Set explicitly rather than relying on build_dir()'s root-relative // default: under BuildOverrides::work_dir the package root is shared // and may be read-only, and the default would write the compiled @@ -8292,6 +8496,9 @@ prepare_build(bool print_fingerprint, // outputs APPENDS to bcRoot.sources, and those appends must be inside // the tail that gets copied into the packages[0] snapshot the scan reads. adoptActionOutputs(*m, *root, ractN); + // The root's build program has spoken; a floor it stated is checked + // now, with the facts every package (it included) established. + if (auto err = checkVersionFloors(); err) return std::unexpected(*err); // Root residues — apply() mutated *m, but packages[0].manifest is a // value-copy snapshot taken at makePackageRoot, so everything the // scan/fingerprint read from the snapshot needs the tail mirrored: @@ -8818,6 +9025,14 @@ prepare_build(bool print_fingerprint, fpi.compileFlags += " #schedule="; fpi.compileFlags += sched; } + // The device axis decides which sources compile and which cfg sections + // apply, so two builds that differ only in it are two builds. Appended + // only when set, so a project that asks for no accelerator keeps the + // build directory it has. + if (const auto accel = resolvedAccel(); !accel.empty()) { + fpi.compileFlags += " #accel="; + fpi.compileFlags += accel; + } if (m->cppStandard.experimental) { // c++fly gate flags are derived (not manifest-declared): fold them in // so a cppfly table change across mcpp versions re-fingerprints. @@ -9044,8 +9259,7 @@ prepare_build(bool print_fingerprint, // no accelerator, and every artifact then satisfies it vacuously — // which is correct, and is why a descriptor lists its CPU-only variant // first: the first accepted artifact wins. - currentTag.accel = mcpp::pack::parse_accel( - overrides.accel.empty() ? m->buildConfig.accel : overrides.accel); + currentTag.accel = mcpp::pack::parse_accel(resolvedAccel()); for (std::size_t i = 1; i < packages.size(); ++i) { auto const& pkg = packages[i]; if (!mcpp::pack::is_distribution_package(pkg.manifest)) continue; @@ -9368,8 +9582,19 @@ prepare_build(bool print_fingerprint, // under plain `mcpp build`, where that unit does not exist. // (`[resources]` makes the opposite call on purpose: an icon // belongs to what ships, not to a test runner.) + // + // ⭐⭐ A STATIC LIBRARY IS ONE OF THEM, and leaving it out was + // the whole of what C-6 needed. A package whose device code is + // its point -- ggml's CUDA backend is 305 `.cu` files behind a + // `kind = "lib"` target -- emitted its actions, watched every + // one of them be dropped with a warning, and produced an + // archive with no device code in it. The archive rule already + // consumes `lu.objects`, so the objects an action produced + // belong there for exactly the reason a compiled `.cpp`'s do: + // the target's content is what it was told to contain. const bool image = lu.kind == mcpp::build::LinkUnit::Binary || lu.kind == mcpp::build::LinkUnit::SharedLibrary + || lu.kind == mcpp::build::LinkUnit::StaticLibrary || lu.kind == mcpp::build::LinkUnit::TestBinary; const bool wanted = a.targets.empty() ? image @@ -9388,11 +9613,12 @@ prepare_build(bool print_fingerprint, if (!attached && a.targets.empty()) { mcpp::diag::degraded("action/no-target", std::format( "build.mcpp action '{}' has role = \"object\" but this build " - "produces no executable, shared library or test binary to " - "link its outputs into", a.id.empty() ? "" : a.id), + "produces no target to put its outputs into", + a.id.empty() ? "" : a.id), "the action never runs and its outputs are never produced", - "add a [targets.] that links, or name the targets " - "explicitly with .target(\"…\")"); + "add a [targets.] — a bin, a lib, a shared lib or a " + "test all take one — or name the targets explicitly with " + ".target(\"…\")"); } } if (!unknownObjectTargets.empty()) { @@ -9406,7 +9632,7 @@ prepare_build(bool print_fingerprint, " targets in this build: [{}]\n" " (a target gated by required_features is absent unless those " "features are active; test binaries exist only under `mcpp " - "test`, so name none and the outputs reach every image " + "test`, so name none and the outputs reach every target " "including them)", bad, known.empty() ? std::string("none") : known)); } diff --git a/src/build/refusal.cppm b/src/build/refusal.cppm index 74a155949..374e0fbe4 100644 --- a/src/build/refusal.cppm +++ b/src/build/refusal.cppm @@ -64,6 +64,16 @@ enum class Code { // provider for a requirement, this one is about two implementations of one // interface being in the same link at all. ExclusiveCapability, + // A package requires more of the machine than the machine was declared to + // have. Distinct from a capability that is missing entirely: here the thing + // exists and is too old. + VersionFloorUnmet, + // A source glob is constrained to a device set this build does not cover. + // Distinct from VersionFloorUnmet (the machine is too old) and from a + // prebuilt artifact's tag mismatch (that refusal is about consuming): here + // the project's own sources ask for a device the build was not told to + // target. + AccelMismatch, Other, // a refusal that has not been given a code yet }; @@ -85,6 +95,8 @@ constexpr std::string_view name(Code c) { case Code::HostToolToolchain: return "host-tool-toolchain"; case Code::StdModulePrecompile: return "std-module-precompile"; case Code::ExclusiveCapability: return "exclusive-capability"; + case Code::VersionFloorUnmet: return "version-floor-unmet"; + case Code::AccelMismatch: return "accel-mismatch"; case Code::Other: return "other"; } return "other"; diff --git a/src/build/version_floor.cppm b/src/build/version_floor.cppm new file mode 100644 index 000000000..c753383c9 --- /dev/null +++ b/src/build/version_floor.cppm @@ -0,0 +1,159 @@ +// mcpp.build.version_floor — "this needs at least that much of something". +// +// WHY THIS EXISTS +// +// Some facts about a machine bound what may be built for it, and the failure +// when they are ignored arrives late. The case this was written for: a device +// runtime must not be newer than the driver it runs against, and when it is, +// the build compiles and links cleanly and then fails at the first allocation +// with a message that names neither the toolkit nor the driver. +// +// Both numbers are knowable before anything is compiled. What mcpp must not do +// is go and ask a vendor's tool for them — `tests/unit/test_runtime_contract` +// refuses provider-specific probes in `src/`, and rightly: an engine that +// learns to run one vendor's probe learns to run four. So the numbers arrive as +// DECLARATIONS and this module compares them. +// +// WHAT THE DECLARATIONS LOOK LIKE +// +// A package that needs something states a floor, in `runtime.requirements`: +// +// [[runtime.requirements]] +// kind = "version-floor" +// value = "cuda.driver >= 12.0" +// +// A package that KNOWS a fact about this machine states it, in +// `runtime.provides`, having established it at install time: +// +// provides = ["cuda.driver=12.4"] +// +// ⚠️ NO VENDOR NAME APPEARS IN THIS FILE, and that is the point rather than a +// coincidence. `cuda.driver` is data flowing through: this module reads a name, +// a relation and a version, and knows nothing about what any of them mean. A +// second backend needs no change here. +// +// WHAT A VERSION IS +// +// A dot-separated sequence of integers, compared component by component, with a +// missing component reading as zero — so `12` and `12.0` are the same version +// and `12.4` is above both. Anything that is not that shape yields no version, +// and an absent version on either side yields no claim: a check that cannot +// reach an answer must not manufacture a refusal. + +export module mcpp.build.version_floor; + +import std; + +export namespace mcpp::build { + +// A parsed ` >= ` requirement. +struct VersionFloor { + std::string name; + std::string version; + bool valid() const { return !name.empty() && !version.empty(); } +}; + +// A parsed `=` fact. +struct VersionFact { + std::string name; + std::string version; + bool valid() const { return !name.empty() && !version.empty(); } +}; + +// Read ` >= `. Whitespace around each part is ignored; any other +// shape yields an invalid result rather than a guess. +VersionFloor parse_version_floor(std::string_view text); + +// Read `=`. The separator is a bare `=` so the spelling matches +// the capability strings a descriptor already writes. +VersionFact parse_version_fact(std::string_view text); + +// Is `have` at or above `want`? Both are dot-separated integers; a missing +// component reads as zero. +// +// Returns std::nullopt when either side is not a version, which callers report +// as "no claim" rather than as a failure. +std::optional version_at_least(std::string_view have, std::string_view want); + +} // namespace mcpp::build + +namespace mcpp::build { + +namespace { + +std::string_view trim(std::string_view s) { + while (!s.empty() && (s.front() == ' ' || s.front() == '\t')) s.remove_prefix(1); + while (!s.empty() && (s.back() == ' ' || s.back() == '\t')) s.remove_suffix(1); + return s; +} + +// The components of a dotted version, or an empty vector when the text is not +// one. A trailing or leading dot makes it not one: `12.` is a typo, not `12`. +std::optional> components(std::string_view v) { + v = trim(v); + if (v.empty()) return std::nullopt; + std::vector out; + long long acc = 0; + bool digits = false; + for (char c : v) { + if (c == '.') { + if (!digits) return std::nullopt; + out.push_back(acc); + acc = 0; + digits = false; + continue; + } + if (!std::isdigit(static_cast(c))) return std::nullopt; + acc = acc * 10 + (c - '0'); + digits = true; + } + if (!digits) return std::nullopt; + out.push_back(acc); + return out; +} + +} // namespace + +VersionFloor parse_version_floor(std::string_view text) { + VersionFloor f; + auto at = text.find(">="); + if (at == std::string_view::npos) return f; + auto name = trim(text.substr(0, at)); + auto ver = trim(text.substr(at + 2)); + if (name.empty() || ver.empty()) return f; + if (!components(ver)) return f; + f.name = std::string(name); + f.version = std::string(ver); + return f; +} + +VersionFact parse_version_fact(std::string_view text) { + VersionFact f; + auto at = text.find('='); + if (at == std::string_view::npos) return f; + // `>=` is a floor, not a fact; refusing it here keeps one spelling from + // being read as the other when both live in string lists. + if (at > 0 && text[at - 1] == '>') return f; + auto name = trim(text.substr(0, at)); + auto ver = trim(text.substr(at + 1)); + if (name.empty() || ver.empty()) return f; + if (!components(ver)) return f; + f.name = std::string(name); + f.version = std::string(ver); + return f; +} + +std::optional version_at_least(std::string_view have, std::string_view want) { + auto h = components(have); + auto w = components(want); + if (!h || !w) return std::nullopt; + const auto n = std::max(h->size(), w->size()); + for (std::size_t i = 0; i < n; ++i) { + const long long a = i < h->size() ? (*h)[i] : 0; + const long long b = i < w->size() ? (*w)[i] : 0; + if (a != b) return a > b; + } + return true; +} + +} // namespace mcpp::build diff --git a/src/cli.cppm b/src/cli.cppm index bd1565dd1..bad95341c 100644 --- a/src/cli.cppm +++ b/src/cli.cppm @@ -409,6 +409,10 @@ int run(int argc, char** argv) { .arg(cl::Arg("bin").help("Binary name (optional)")) .option(cl::Option("target").takes_value().value_name("TRIPLE") .help("Cross target triple (same axis as `mcpp build --target`)")) + .option(cl::Option("accel").takes_value().value_name("SPEC") + .help("Device backends and architectures (same axis as `mcpp build --accel`)")) + .option(cl::Option("no-accel") + .help("Run the variant built for no accelerator (same as `mcpp build --no-accel`)")) // Kept as an alias: it shipped in 2026.8.19.1 as the only spelling // `run` accepted, and scripts written against it must keep working. .option(cl::Option("target-triple").takes_value().value_name("TRIPLE") @@ -466,6 +470,10 @@ int run(int argc, char** argv) { .help("Run only tests whose name contains PATTERN (optional)")) .option(cl::Option("target").takes_value().value_name("TRIPLE") .help("Cross target triple (same axis as `mcpp build --target`)")) + .option(cl::Option("accel").takes_value().value_name("SPEC") + .help("Device backends and architectures (same axis as `mcpp build --accel`)")) + .option(cl::Option("no-accel") + .help("Test the variant built for no accelerator (same as `mcpp build --no-accel`)")) .option(cl::Option("message-format").takes_value().value_name("FMT") .help("Output format: human (default) | json (NDJSON, one record per test)")) .option(cl::Option("list") diff --git a/src/cli/cmd_build.cppm b/src/cli/cmd_build.cppm index 18a476cbb..c04484863 100644 --- a/src/cli/cmd_build.cppm +++ b/src/cli/cmd_build.cppm @@ -178,10 +178,16 @@ export int cmd_build(const mcpplibs::cmdline::ParsedArgs& parsed) { // override (--profile/--features/--strict, like --target/--static) must // bypass it: the cached build.ninja was generated without them, so taking // the fast path would silently ignore the flags. + // + // `--accel` / `--no-accel` are on that list. Measured before they were: a + // successful `mcpp build` followed by `mcpp build --no-accel` reported + // "Finished in 0.00s" and handed back the device build -- the axis that + // decides which sources compile and which cfg sections apply, ignored + // because the flag arrived after a build that did not carry it. if (!print_fp && ov.target_triple.empty() && !ov.force_static && ov.profile.empty() && ov.features.empty() && !ov.strict && ov.capabilities.empty() && ov.package_filter.empty() - && ov.cache_mode.empty()) { + && ov.cache_mode.empty() && ov.accel.empty()) { auto root = mcpp::project::find_manifest_root(std::filesystem::current_path()); if (root) { // A project with active `[hooks]` declines the fast path from @@ -235,13 +241,20 @@ export int cmd_run(const mcpplibs::cmdline::ParsedArgs& parsed, if (auto pr = parsed.value("profile")) profile = *pr; if (parsed.is_flag_set("release")) profile = "release"; if (parsed.is_flag_set("dev")) profile = "dev"; + // The device axis, read exactly as `build` reads it: `--no-accel` is an + // explicit choice and not the absence of `--accel`, so it travels as the + // same sentinel. Without this a project's CPU-only variant could be built + // but not run through the command surface. + std::string accel; + if (parsed.is_flag_set("no-accel")) accel = "(none)"; + else if (auto a = parsed.value("accel")) accel = *a; if (parsed.is_flag_set("list-runners")) return mcpp::build::list_runners(package_filter, cache_mode, no_cache, - target_triple, features, profile); + target_triple, features, profile, accel); return mcpp::build::build_run_target(targetName, passthrough, package_filter, cache_mode, no_cache, target_triple, no_runner, runner_name, features, - profile); + profile, accel); } export int cmd_test(const mcpplibs::cmdline::ParsedArgs& parsed, @@ -261,6 +274,8 @@ export int cmd_test(const mcpplibs::cmdline::ParsedArgs& parsed, else if (parsed.is_flag_set("no-cache")) ov.cache_mode = "off"; if (auto tt = parsed.value("target")) ov.target_triple = *tt; + if (parsed.is_flag_set("no-accel")) ov.accel = "(none)"; + else if (auto a = parsed.value("accel")) ov.accel = *a; mcpp::build::TestOptions to; if (parsed.positional_count() > 0) to.filter = parsed.positional(0); diff --git a/src/config.cppm b/src/config.cppm index cd49375ae..d26ab064e 100644 --- a/src/config.cppm +++ b/src/config.cppm @@ -625,13 +625,40 @@ std::expected load_or_init( // TODO(xlings-upstream): collapse into a single // `xlings sandbox bootstrap --home ` once that command exists // upstream (see docs/short-term-vs-long-track plan). + // + // NOT UNDER --offline. Each of the three steps below reaches the + // network when the sandbox is fresh: `xlings self init` clones the + // package index, and patchelf and ninja arrive through `xlings + // install`. `--offline` promises never to touch the network, and an + // empty home is exactly where that promise was being broken -- + // measured: `MCPP_OFFLINE=1 mcpp self doctor` in a fresh home fetched + // the index and 126 MB of tools before running its first check. + // + // A bootstrapped home loses nothing: every step is a no-op once its + // artefact exists, and the completion markers below are still written + // from what is on disk. A fresh home stays un-bootstrapped, and the + // commands that need the tools say so through check_base_init(). + const bool bootstrapAllowed = !mcpp::platform::env::offline_mode(); ensure_sandbox_xlings_binary(cfg, quiet); - ensure_sandbox_init(cfg, quiet); + if (bootstrapAllowed) { + ensure_sandbox_init(cfg, quiet); + } else if (!quiet) { + // Once per process: the doctor loads the configuration more than once, + // and the same sentence twice reads as two events. + static bool announced = false; + auto marker = mcpp::xlings::paths::sandbox_init_marker(make_xlings_env(cfg)); + if (!announced && !std::filesystem::exists(marker)) { + announced = true; + print_status("Skipping", + "sandbox bootstrap (offline mode; run without --offline to bootstrap)"); + } + } { auto bsEnv = make_xlings_env(cfg); #if !defined(__APPLE__) && !defined(_WIN32) // patchelf is ELF-only; macOS uses Mach-O and Windows uses PE. - ensure_sandbox_patchelf(cfg, quiet, onBootstrapProgress); + if (bootstrapAllowed) + ensure_sandbox_patchelf(cfg, quiet, onBootstrapProgress); // Only mark complete if the actual binary exists (not just the dir). { auto pBin = mcpp::xlings::paths::xim_tool(bsEnv, "patchelf", @@ -640,7 +667,8 @@ std::expected load_or_init( mcpp::fallback::mark_install_complete(pBin.parent_path().parent_path()); } #endif - ensure_sandbox_ninja(cfg, quiet, onBootstrapProgress); + if (bootstrapAllowed) + ensure_sandbox_ninja(cfg, quiet, onBootstrapProgress); { auto nRoot = mcpp::xlings::paths::xim_tool_root(bsEnv, "ninja"); auto ninja_name = std::string("ninja") + std::string(mcpp::platform::exe_suffix); diff --git a/src/doctor.cppm b/src/doctor.cppm index 27fbe5ace..b8f12a996 100644 --- a/src/doctor.cppm +++ b/src/doctor.cppm @@ -34,7 +34,6 @@ import mcpp.runtime.elf; import mcpp.pm.index_refresh; // staleness_note for `mcpp why deps` import mcpp.project; import mcpp.toolchain.detect; -import mcpp.toolchain.devicehost; import mcpp.toolchain.msvc; import mcpp.toolchain.registry; import mcpp.toolchain.linkmodel; @@ -102,69 +101,6 @@ export int env_report() { } // `mcpp self doctor`. -// Which back-end stage nvcc names but cannot resolve, if any. -// -// Nothing is compiled: `--dryrun` prints the plan and stops. std::nullopt -// covers three unlike situations on purpose -- there is no nvcc, the dryrun -// produced no plan, and every stage in the plan resolves -- because only a -// named unresolvable stage is a finding. A probe that cannot reach an answer -// must not manufacture one. -std::optional unreachable_device_stage() { - namespace fs = std::filesystem; - std::error_code ec; - - // A fresh directory per run, on the pattern the p1689 scanner already - // uses. A fixed name under the shared temporary directory would be a - // path another user can create first, and the `remove_all` that a fixed - // name needs in order to be reusable is the part that makes that matter. - const auto probe = fs::temp_directory_path(ec) - / std::format("mcpp_nvcc_dryrun_{}", std::random_device{}()); - if (ec) return std::nullopt; - // The return value, not `ec`: create_directory reports an existing - // directory by returning false without setting an error, and proceeding - // into a directory this process did not create is the case being avoided. - if (!fs::create_directory(probe, ec) || ec) return std::nullopt; - struct Cleanup { - fs::path dir; - ~Cleanup() { std::error_code e; fs::remove_all(dir, e); } - } const cleanup{probe}; - - const auto source = probe / "empty.cu"; - { std::ofstream out(source); if (!out) return std::nullopt; } - - // A spawn that fails because there is no nvcc yields text with no `#$` - // lines, hence an empty plan, hence no finding. No separate check needed. - const auto run = mcpp::platform::process::capture_exec( - {"nvcc", "--dryrun", "-c", source.string(), - "-o", (probe / "empty.o").string()}); - - const auto plan = mcpp::toolchain::parse_dryrun(run.output); - if (plan.programs.empty()) return std::nullopt; - - std::string search = plan.searchPath; - if (search.empty()) - if (const char* p = std::getenv("PATH"); p) search = p; - if (search.empty()) return std::nullopt; - - std::vector dirs; - for (std::size_t pos = 0; pos <= search.size(); ) { - const auto sep = search.find(':', pos); - const auto stop = sep == std::string::npos ? search.size() : sep; - if (stop > pos) dirs.emplace_back(search.substr(pos, stop - pos)); - pos = stop + 1; - } - - for (auto const& program : plan.programs) { - bool found = false; - for (auto const& dir : dirs) { - if (fs::exists(dir / program, ec)) { found = true; break; } - ec.clear(); - } - if (!found) return program; - } - return std::nullopt; -} - export int doctor_report() { int warns = 0, errors = 0; auto ok = [](std::string_view m) { mcpp::ui::status("ok", m); }; @@ -641,108 +577,15 @@ export int doctor_report() { ok("process deadlines: enforced (POSIX SIGKILL / Windows job object)"); } - // ── Device toolkit, and whether this host compiler can drive it ──────── - // - // WHY THIS IS A DOCTOR CHECK AND NOT A BUILD ERROR - // - // nvcc refuses host compilers newer than a bound the toolkit states in its - // own crt/host_config.h. The failure is late, the message names a compiler - // the user did not choose, and the bound is invisible. Every other build - // system forwards -ccbin and lets nvcc discover this; mcpp supplies the - // host compiler and can therefore answer before anything is compiled. - // - // Reported rather than enforced because a project that compiles no device - // code is unaffected by an incompatible pair, and refusing its build would - // be a false alarm. - if (!mcpp::platform::is_windows) { - auto header = [&]() -> std::optional { - std::vector roots; - if (const char* p = std::getenv("CUDA_PATH")) roots.emplace_back(p); - if (const char* p = std::getenv("CUDA_HOME")) roots.emplace_back(p); - roots.emplace_back("/usr/local/cuda"); - std::error_code ec; - for (auto const& r : roots) { - auto h = r / "include" / "crt" / "host_config.h"; - if (std::filesystem::exists(h, ec)) return h; - } - // Distribution packaging puts the toolkit headers on the default - // include path instead of under a versioned root. - std::filesystem::path sys = "/usr/include/crt/host_config.h"; - if (std::filesystem::exists(sys, ec)) return sys; - return std::nullopt; - }(); - - if (header) { - mcpp::ui::status("Checking", "device toolkit"); - std::ifstream in(*header); - std::string text((std::istreambuf_iterator(in)), - std::istreambuf_iterator()); - auto bounds = mcpp::toolchain::parse_host_config(text); - if (!bounds.known()) { - ok(std::format("cuda headers at {} (no host-compiler bound stated)", - header->parent_path().parent_path().string())); - } else if (!tc) { - ok(std::format("cuda host-compiler bound: gcc<={} clang<={}", - bounds.gccMax, bounds.clangMax)); - } else { - const std::string family(tc->compiler_name()); - // Leading digits of the version. Extracted here rather than - // reached for from mcpp.pack.abi_tag: three lines are cheaper - // than a module edge from the diagnostics layer to packaging. - const int major = [&] { - int v = 0; - for (char c : tc->version) { - if (!std::isdigit(static_cast(c))) break; - v = v * 10 + (c - '0'); - } - return v; - }(); - if (mcpp::toolchain::host_compiler_accepted(bounds, family, major)) { - ok(std::format("cuda accepts this host compiler ({} {} <= {})", - family, major, - family == "gcc" ? bounds.gccMax : bounds.clangMax)); - } else { - warn(std::format( - "cuda will refuse this host compiler: {} {} exceeds the " - "bound of {} stated in {}.\n" - " Device code will not compile until a host " - "compiler within the bound is selected; a project that " - "compiles no device code is unaffected.", - family, major, - family == "gcc" ? bounds.gccMax : bounds.clangMax, - header->string())); - } - } - - // WHETHER nvcc CAN REACH ITS OWN BACK-END - // - // A toolkit can be present, complete and on PATH and still fail - // at the first stage, because nvcc resolves cicc, ptxas and - // fatbinary as bare names on a PATH it prepends from an - // `nvcc.profile` beside its binary. A container or sandbox that - // replaces /etc removes that profile -- it is a symlink into it - // on Debian-family packaging -- and nvcc then states no PATH and - // reports `sh: 1: cicc: not found`. The message names neither - // nvcc nor the profile, and nothing about the toolkit is missing, - // so the user has nowhere to look. - // - // Asked rather than assumed: `--dryrun` prints the plan without - // running it, so the answer is nvcc's own. - if (auto missing = unreachable_device_stage(); missing) { - warn(std::format( - "nvcc cannot reach its own back-end: it invokes '{}' by " - "name, and that name does not resolve on the search path " - "it states.\n" - " The toolkit is installed; what is missing is the " - "`nvcc.profile` that prepends the toolkit's own bin " - "directory. This is what a container or sandbox that " - "replaces /etc removes. Device code will fail to compile " - "with a message naming only '{}'.", - *missing, *missing)); - } - } - } - + // No device-toolkit report here, on purpose. The bound a toolkit states + // for its host compiler, whether a device compiler can reach its own + // back-end, and whether the driver is new enough for the runtime are all + // facts about one vendor's tools, and the rule package that drives those + // tools is where they are read: it reports through `mcpp::warning` and + // states the driver relation through `mcpp::fact` / `mcpp::floor`, and the + // engine compares before the first compile. A machine with no rule + // package in its project has nothing vendor-specific to say here, and says + // nothing (`tests/unit/test_core_vendor_probes.cpp` holds that line). std::println(""); if (errors) std::println("Doctor result: {} errors, {} warnings", errors, warns); else if (warns) std::println("Doctor result: {} warnings", warns); diff --git a/src/toolchain/devicehost.cppm b/src/toolchain/devicehost.cppm deleted file mode 100644 index ac2a74d1e..000000000 --- a/src/toolchain/devicehost.cppm +++ /dev/null @@ -1,174 +0,0 @@ -// mcpp.toolchain.devicehost — which host compilers a device toolchain accepts. -// -// WHY THIS IS READ RATHER THAN TABULATED -// -// nvcc refuses host compilers newer than a bound that changes with every CUDA -// release, and the bound is not documentation: it is a preprocessor guard in -// the toolkit's own `crt/host_config.h`. A table transcribed into mcpp would -// be a copy of one release, correct until the next one and wrong silently -// afterwards, and it would have to grow a row for every future toolkit. -// Reading the guard means a toolkit mcpp has never heard of still answers. -// -// WHAT IT IS FOR -// -// mcpp supplies the host compiler, so it is the one build system in a position -// to know both sides of this pairing before either compiler runs. CMake -// forwards `-ccbin` and lets nvcc fail; the diagnostic then names a compiler -// the user did not choose and a bound they cannot see. Refusing earlier, with -// both versions and the bound in the message, is the whole of the benefit. -// -// The parse is deliberately narrow: two guards, no macro evaluation. A file -// this module cannot make sense of yields no bound, and no bound means the -// check does not run — an unreadable header must not invent a refusal. - -export module mcpp.toolchain.devicehost; - -import std; - -export namespace mcpp::toolchain { - -// The host-compiler bounds one device toolkit declares. Zero means "the header -// said nothing about this family", never "no version is allowed". -struct HostCompilerBounds { - int gccMax = 0; // greatest accepted __GNUC__ - int clangMax = 0; // greatest accepted clang major - - bool known() const { return gccMax != 0 || clangMax != 0; } -}; - -// Parse the two guards out of a `crt/host_config.h`. -HostCompilerBounds parse_host_config(std::string_view header); - -// The plan nvcc states for one compilation: the search path it will use, and -// the programs it will invoke by bare name. -// -// WHY THE PLAN IS ASKED FOR RATHER THAN ASSUMED -// -// nvcc runs its back-end stages -- cicc, ptxas, fatbinary, nvlink -- as bare -// names, on a PATH it prepends itself from an `nvcc.profile` beside its own -// binary. Where those stages live is therefore not a property of the toolkit -// layout that mcpp could tabulate: it is whatever that profile says. When the -// profile is unreachable -- a container or sandbox that replaces /etc, where -// the profile is a symlink into it -- nvcc states no PATH, keeps the ambient -// one, and fails at the first stage with `sh: 1: cicc: not found`. That -// message names neither nvcc nor the profile, and the toolkit is present and -// intact, so every obvious check passes. `--dryrun` prints the same plan -// without running it, which is how the question is answered before a -// compilation is attempted. -struct DeviceDryRun { - // The PATH nvcc assigns itself. Empty means it assigned none, in which - // case the ambient PATH is what its stages will be resolved against. - std::string searchPath; - // Stages invoked by bare name, in first-appearance order, deduplicated. - // A stage named by an absolute or relative path resolves on its own and - // is not collected. - std::vector programs; -}; - -// Parse the `#$` lines of `nvcc --dryrun` output. -DeviceDryRun parse_dryrun(std::string_view text); - -// Is `major` of `family` ("gcc" | "clang") within the bounds? A family the -// header said nothing about is accepted: silence is not a refusal. -bool host_compiler_accepted(const HostCompilerBounds& b, - std::string_view family, int major); - -} // namespace mcpp::toolchain - -namespace mcpp::toolchain { - -namespace { - -// The first run of digits at or after `pos`, or 0. -int digits_after(std::string_view s, std::size_t pos) { - while (pos < s.size() && !std::isdigit(static_cast(s[pos]))) { - // Stop at a line break: a number on the next line belongs to another - // statement, and reading across one is how a parse this narrow would - // start inventing answers. - if (s[pos] == '\n') return 0; - ++pos; - } - int v = 0, n = 0; - while (pos < s.size() && std::isdigit(static_cast(s[pos]))) { - v = v * 10 + (s[pos] - '0'); - ++pos; ++n; - } - return n == 0 ? 0 : v; -} - -} // namespace - -HostCompilerBounds parse_host_config(std::string_view header) { - HostCompilerBounds b; - - // `#if __GNUC__ > 12` guards the "gcc versions later than 12" error, so - // the greatest accepted major is the operand itself. - if (auto p = header.find("__GNUC__ > "); p != std::string_view::npos) - b.gccMax = digits_after(header, p + 10); - - // The clang guard states its bound in prose rather than in the condition: - // "clang version must be less than 15 and greater than 3.2". - if (auto p = header.find("clang version must be less than "); - p != std::string_view::npos) { - int exclusive = digits_after(header, p + 31); - if (exclusive > 0) b.clangMax = exclusive - 1; - } - return b; -} - -DeviceDryRun parse_dryrun(std::string_view text) { - DeviceDryRun plan; - - for (std::size_t pos = 0; pos <= text.size(); ) { - const auto eol = text.find('\n', pos); - const auto stop = eol == std::string_view::npos ? text.size() : eol; - std::string_view line = text.substr(pos, stop - pos); - pos = stop + 1; // past the end after the last line: the loop stops - - // Every line nvcc contributes is prefixed; anything else is a - // diagnostic and says nothing about the plan. - constexpr std::string_view kPrefix = "#$ "; - if (!line.starts_with(kPrefix)) continue; - line.remove_prefix(kPrefix.size()); - while (!line.empty() && (line.front() == ' ' || line.front() == '\t')) - line.remove_prefix(1); - while (!line.empty() && (line.back() == '\r' || line.back() == ' ')) - line.remove_suffix(1); - if (line.empty()) continue; - - // The first token ends at whitespace or at the `=` of an assignment, - // whichever comes first. An assignment is `NAME=value` with no space - // before the `=`; a command is anything else. - std::size_t end = 0; - while (end < line.size() && line[end] != ' ' && line[end] != '\t' - && line[end] != '=') - ++end; - - if (end < line.size() && line[end] == '=') { - if (line.substr(0, end) == "PATH") - plan.searchPath = std::string(line.substr(end + 1)); - continue; - } - - std::string_view program = line.substr(0, end); - // A stage nvcc names by path resolves without the search path. - if (program.find('/') != std::string_view::npos - || program.find('\\') != std::string_view::npos - || program.starts_with("\"")) continue; - if (std::ranges::find(plan.programs, program) == plan.programs.end()) - plan.programs.emplace_back(program); - } - return plan; -} - -bool host_compiler_accepted(const HostCompilerBounds& b, - std::string_view family, int major) -{ - if (major <= 0) return true; // unknown version: no claim - if (family == "gcc") return b.gccMax == 0 || major <= b.gccMax; - if (family == "clang" || family == "llvm") - return b.clangMax == 0 || major <= b.clangMax; - return true; -} - -} // namespace mcpp::toolchain diff --git a/src/toolchain/gcc.cppm b/src/toolchain/gcc.cppm index ceeabbab9..d50b4c751 100644 --- a/src/toolchain/gcc.cppm +++ b/src/toolchain/gcc.cppm @@ -22,6 +22,18 @@ void enrich_toolchain(Toolchain& tc); std::optional find_binutils_bin(const std::filesystem::path& compilerBin); +// The external GNU binutils directory this toolchain must be pointed at with +// `-B`, or empty when it must not be pointed at one — which is every toolchain +// that is not a glibc GCC. +// +// ⭐ ONE STATEMENT OF A GUARD THAT HAD THREE COPIES. musl-cross-make and +// MinGW-w64 payloads bundle their own as/ld, and for a cross target the host's +// binutils would mis-assemble — the Linux `as` rejects MinGW's PE/SEH +// directives `.def` / `.seh_proc`. Only the glibc GCC needs the external +// package. Each copy of that sentence was a place the next reader had to +// re-derive it, and one of them said so in a comment. +std::filesystem::path binutils_prefix_dir(const Toolchain& tc); + std::filesystem::path std_bmi_path(const std::filesystem::path& cacheDir); std::filesystem::path staged_std_bmi_path(const std::filesystem::path& outputDir); @@ -148,6 +160,17 @@ find_binutils_bin(const std::filesystem::path& compilerBin) { return std::nullopt; } +std::filesystem::path binutils_prefix_dir(const Toolchain& tc) { + // GCC only, and the restriction is what makes the name true. Clang resolves + // its assembler and linker through its own payload and must never be handed + // a GNU binutils directory; the engine's clang command lines carry no `-B` + // at all, so answering with one would describe a flag nobody passes. + if (tc.compiler != CompilerId::GCC) return {}; + if (is_musl_target(tc) || is_mingw_target(tc)) return {}; + if (auto bin = find_binutils_bin(tc.binaryPath)) return *bin; + return {}; +} + std::filesystem::path std_bmi_path(const std::filesystem::path& cacheDir) { return cacheDir / "gcm.cache" / "std.gcm"; } @@ -160,17 +183,9 @@ std::string std_module_build_command(const Toolchain& tc, const std::filesystem::path& cacheDir, std::string_view sysrootFlag, std::string_view cppStandardFlag) { - // musl-cross-make AND MinGW-w64 cross toolchains bundle their own as/ld - // (and for a cross target the host's binutils would mis-assemble — e.g. - // the Linux `as` chokes on MinGW's PE/SEH directives `.def`/`.seh_proc`). - // Only the glibc gcc needs an external binutils package wired via -B. - // Mirrors the guard in build/flags.cppm. std::string bFlag; - if (!is_musl_target(tc) && !is_mingw_target(tc)) { - if (auto binutilsBin = find_binutils_bin(tc.binaryPath)) { - bFlag = std::format(" -B{}", mcpp::xlings::shq(binutilsBin->string())); - } - } + if (auto binutilsBin = binutils_prefix_dir(tc); !binutilsBin.empty()) + bFlag = std::format(" -B{}", mcpp::xlings::shq(binutilsBin.string())); // Windows (MinGW): cmd.exe needs `/d` to change DRIVE (project on D:, // BMI cache on C: is the real CI layout — same-drive runs masked this). diff --git a/tests/e2e/188_build_actions.sh b/tests/e2e/188_build_actions.sh index 1d88f0c46..ded749002 100755 --- a/tests/e2e/188_build_actions.sh +++ b/tests/e2e/188_build_actions.sh @@ -413,6 +413,13 @@ grep -q '1 passed' o6.log || { cat o6.log; echo "FAIL: the test did not run"; ex # supposed to be reachable through a link edge), so "no consumer" means the # command never runs. Saying nothing there is the same failure `[resources]` # reports as resources/no-image. +# +# ⚠️ THE ONLY TARGET IS GATED OFF, and that is what makes this case exist at +# all. A `kind = "lib"` target used to serve here, because a static library did +# not take an object action's outputs; it does now (design C-6, e2e 608), so a +# project with one is no longer a project with no consumer. A target behind an +# inactive `required_features` is absent from the build entirely, which is the +# one shape that leaves an object action with nowhere to go. mkdir -p "$TMP/objnone/src" cd "$TMP/objnone" cat > mcpp.toml <<'EOF' @@ -420,17 +427,30 @@ cat > mcpp.toml <<'EOF' name = "objnone" version = "0.1.0" +[features] +gated = {} + [targets.objnone] kind = "lib" +required_features = ["gated"] EOF printf 'export module objnone;\nexport int f(){return 1;}\n' > src/objnone.cppm cp "$TMP/objtest/blob.cpp" blob.cpp cp "$TMP/objtest/mkobj.sh" mkobj.sh cp "$TMP/objtest/build.mcpp" build.mcpp "$MCPP" build > o7.log 2>&1 || { cat o7.log; echo "FAIL: objnone build failed"; exit 1; } -grep -q 'no executable, shared library or test binary' o7.log || { +grep -q 'produces no target to put its outputs into' o7.log || { cat o7.log; echo "FAIL: an object with no consumer must be reported"; exit 1; } +# ⭐ THE CONTROL: with the feature on, the target exists and the action is not +# reported. Without it the assertion above would pass on a build that had +# stopped emitting the diagnostic for every project. +"$MCPP" build --features gated > o7b.log 2>&1 || { + cat o7b.log; echo "FAIL: objnone build with the feature failed"; exit 1; } +grep -q 'produces no target to put its outputs into' o7b.log && { + cat o7b.log; echo "FAIL: the object was reported as having no consumer while its target existed" + exit 1; } + cd "$TMP/edge" # ── 4. a malformed action is refused, not skipped ────────────────────────── diff --git a/tests/e2e/317_project_build_hooks.sh b/tests/e2e/317_project_build_hooks.sh index 0b4e6613d..c32972f23 100644 --- a/tests/e2e/317_project_build_hooks.sh +++ b/tests/e2e/317_project_build_hooks.sh @@ -182,9 +182,11 @@ esac if [[ $HOST_WINDOWS -eq 1 ]]; then slow_command="ping -n 6 127.0.0.1 >NUL" pause_2s="ping -n 3 127.0.0.1 >NUL" + pause_5s="ping -n 6 127.0.0.1 >NUL" else slow_command="sleep 5" pause_2s="sleep 2" + pause_5s="sleep 5" fi write_manifest < "$1/mcpp.toml" + echo "int ${1//-/_}_anchor(void){return 0;}" > "$1/src/$1.c" +} + +# A package that needs something of the machine. +mkneed() { # $1 name $2 "cuda.driver >= 13.0" + mkdir -p "$1/src" + { + printf '[package]\nname = "%s"\nversion = "0.1.0"\n[language]\nstandard = "c++23"\n' "$1" + printf '[[runtime.requirements]]\nkind = "version-floor"\nvalue = "%s"\n' "$2" + printf '[targets.%s]\nkind = "lib"\n' "$1" + } > "$1/mcpp.toml" + echo "int ${1//-/_}_anchor(void){return 0;}" > "$1/src/$1.c" +} + +consumer() { + rm -rf app; "$MCPP" new app > /dev/null; cd app + { printf '[package]\nname="app"\nversion="0.1.0"\n[language]\nstandard="c++23"\n[dependencies]\n' + for d in "$@"; do echo "$d = { path = \"../$d\" }"; done; } > mcpp.toml + rm -f src/*.cppm + echo 'int main(){return 0;}' > src/main.cpp +} + +# ── One: the floor is above the fact ─────────────────────────────────────── +mkfact driverfact "cuda.driver=12.4" +mkneed toolkitnew "cuda.driver >= 13.0" +consumer driverfact toolkitnew + +if "$MCPP" build > out.log 2>&1; then + cat out.log; echo "FAIL: a floor above the stated fact was accepted"; exit 1 +fi +grep -q "cuda.driver" out.log || { cat out.log; echo "FAIL: refusal does not name what is short"; exit 1; } +grep -q "13.0" out.log || { cat out.log; echo "FAIL: refusal does not say what was needed"; exit 1; } +grep -q "12.4" out.log || { cat out.log; echo "FAIL: refusal does not say what is there"; exit 1; } +grep -q "driverfact" out.log || { cat out.log; echo "FAIL: refusal does not say who stated the fact"; exit 1; } +echo "PASS: refused, naming the requirement, both versions and who stated the fact" + +# ── Two: the floor is met ────────────────────────────────────────────────── +cd "$TMP" +mkneed toolkitok "cuda.driver >= 12.0" +consumer driverfact toolkitok +"$MCPP" build > ok.log 2>&1 || { cat ok.log; echo "FAIL: a met floor was refused"; exit 1; } +echo "PASS: a met floor builds" + +# ── Three: a floor with NO fact is silent, and the name is just data ─────── +# +# This is the control that matters. Without it the first case would also pass +# against an engine that refused every version-floor requirement, which would +# turn "we do not know" into "no" — the exact failure mode this mechanism was +# built to avoid. +cd "$TMP" +mkneed futureneed "some.future.thing >= 4.2.1" +consumer futureneed +"$MCPP" build > quiet.log 2>&1 || { cat quiet.log; echo "FAIL: a floor nobody answered was refused"; exit 1; } +echo "PASS: a floor with no stated fact is silent, for a name no backend uses" + +echo "PASS: version floor" diff --git a/tests/e2e/604_offline_skips_sandbox_bootstrap.sh b/tests/e2e/604_offline_skips_sandbox_bootstrap.sh new file mode 100755 index 000000000..9a8fc5410 --- /dev/null +++ b/tests/e2e/604_offline_skips_sandbox_bootstrap.sh @@ -0,0 +1,49 @@ +#!/usr/bin/env bash +# `--offline` promises never to touch the network, and a fresh home is where +# that promise was being broken: the sandbox bootstrap that `load_or_init` runs +# on first use clones the package index and installs ninja and patchelf through +# `xlings install`. Measured before the gate, `MCPP_OFFLINE=1 mcpp self doctor` +# in an empty home spent 26 seconds and wrote 126 MB before its first check. +# +# Under offline mode the three network-bound steps are skipped, once, visibly, +# and the command goes on to its remaining checks. A home that is already +# bootstrapped says nothing, because there is nothing to skip. +set -e + +TMP=$(mktemp -d) +trap "rm -rf $TMP" EXIT + +# ── A fresh home under offline mode is left un-bootstrapped ───────────── +out="$TMP/fresh.log" +MCPP_HOME="$TMP/home" MCPP_OFFLINE=1 "$MCPP" self doctor > "$out" 2>&1 || true + +n=$(grep -c "Skipping sandbox bootstrap (offline mode" "$out" || true) +[[ "$n" -eq 1 ]] || { + cat "$out"; echo "FAIL: expected the skip to be announced once, saw $n"; exit 1; } + +# The index clone is the first network step and the largest; its absence is +# the structural evidence that nothing was fetched. +[[ ! -e "$TMP/home/registry/data/xim-pkgindex" ]] || { + echo "FAIL: the package index was cloned under offline mode"; exit 1; } +[[ ! -e "$TMP/home/registry/subos/default/.xlings.json" ]] || { + echo "FAIL: the sandbox was initialised under offline mode"; exit 1; } + +# Skipping the bootstrap must not end the command: the checks after the +# registry one still run. +grep -q "Checking build policy" "$out" || { + cat "$out"; echo "FAIL: the doctor stopped at the registry check"; exit 1; } +echo "PASS: a fresh home under offline mode is left un-bootstrapped, audibly" + +# ── Control: a bootstrapped home has nothing to skip and says nothing ─── +# +# Without this, a doctor that printed the sentence on every offline run would +# pass the count above. The e2e runner's own home is bootstrapped by the time +# this script runs. +out2="$TMP/warm.log" +MCPP_OFFLINE=1 "$MCPP" self doctor > "$out2" 2>&1 || true +if grep -q "Skipping sandbox bootstrap" "$out2"; then + cat "$out2"; echo "FAIL: a bootstrapped home announced a skip"; exit 1 +fi +echo "PASS: a bootstrapped home under offline mode announces nothing" + +echo "PASS: offline skips the sandbox bootstrap" diff --git a/tests/e2e/605_probe_channel_and_accel_reach_the_build_program.sh b/tests/e2e/605_probe_channel_and_accel_reach_the_build_program.sh new file mode 100755 index 000000000..1270c1977 --- /dev/null +++ b/tests/e2e/605_probe_channel_and_accel_reach_the_build_program.sh @@ -0,0 +1,124 @@ +#!/usr/bin/env bash +# requires: gcc +# The probe channel, and the device axis as a build program sees it. +# +# A rule package is the thing that knows how to ask a machine what it has, and +# the engine is the thing that must not. So the package MEASURES and the engine +# COMPARES: `mcpp::fact` states what the machine has, `mcpp::floor` what the +# package needs of it, and the build is refused before anything is compiled +# when the floor is unmet -- naming both values; `mcpp why toolchain --format +# json` classifies the outcome under the reason +# `version-floor-unmet`. The root's build program is where a rule package +# speaks from, so the check has to run AFTER it; before this test it ran only +# before, and a floor stated there was never compared. +# +# The same program reads the resolved `accel` (`MCPP_ACCEL`), and the +# `cfg(accelerator = "...")` layer key is fed from the same value -- a key that +# was declared, documented, and never written before this. +# +# Nothing here names a vendor: `widget` is the backend, and the engine treats +# it exactly as it would any other. That is the property. +set -e + +TMP=$(mktemp -d) +trap "rm -rf $TMP" EXIT +cd "$TMP" + +"$MCPP" new probe > /dev/null; cd probe +rm -f src/*.cppm +cat > src/main.cpp <<'EOF' +// Compiles only when the accelerator layer selected the widget backend: the +// define arrives through a cfg section keyed on `accelerator`. +#ifndef WIDGET_ON +#error "cfg(accelerator = \"widget\") did not match, so WIDGET_ON is missing" +#endif +int main() { return 0; } +EOF + +write_manifest() { # $1 = floor spec + cat > mcpp.toml < build.mcpp <= 2.0" +if "$MCPP" build > refused.log 2>&1; then + cat refused.log; echo "FAIL: a floor above the stated fact was accepted"; exit 1 +fi +grep -q "widget.driver" refused.log || { cat refused.log; echo "FAIL: refusal does not name the quantity"; exit 1; } +grep -q "2.0" refused.log || { cat refused.log; echo "FAIL: refusal does not say what was needed"; exit 1; } +grep -q "1.2" refused.log || { cat refused.log; echo "FAIL: refusal does not say what is there"; exit 1; } +echo "PASS: a floor stated by the root's build program is compared, and refused with both values" + +# The machine interface names the reason, so a tool can act on it without +# parsing prose. `why toolchain` runs the same prepare, build program included. +"$MCPP" why toolchain --format json > refused.json 2>/dev/null || true +reason="$(jq -r '.data.reason // "-"' refused.json | tr -d '\r')" +[[ "$reason" == "version-floor-unmet" ]] || { + cat refused.json; echo "FAIL: reason is '$reason', expected version-floor-unmet"; exit 1; } +echo "PASS: the reason is version-floor-unmet under why toolchain --format json" + +# ── Two: the control -- a met floor builds ────────────────────────────── +write_manifest "widget.driver >= 1.0" +touch src/main.cpp +"$MCPP" build > ok.log 2>&1 || { cat ok.log; echo "FAIL: a met floor was refused"; exit 1; } +echo "PASS: a met floor builds" + +# ── Three: the build program saw the resolved accel ─────────────────────── +# +# The advisory carries whatever `mcpp::accel()` returned. Parsed and printed +# back by the engine, so the spelling is canonical whatever the manifest wrote. +grep -q "accel=\[widget9+{w1,w2}\]" ok.log || { + cat ok.log; echo "FAIL: MCPP_ACCEL did not carry the resolved accel"; exit 1; } +echo "PASS: MCPP_ACCEL carries the resolved accel" + +# ── Four: --no-accel empties both the variable and the layer ────────────── +# +# The layer is the sharper half: with no backend enabled the cfg section must +# NOT apply, and then main.cpp's #error fires. A key that still matched here +# would be reading the manifest rather than the build. +if "$MCPP" build --no-accel > none.log 2>&1; then + cat none.log; echo "FAIL: cfg(accelerator) matched under --no-accel"; exit 1 +fi +grep -q "WIDGET_ON is missing" none.log || { + cat none.log; echo "FAIL: the build failed for another reason than the layer"; exit 1; } +echo "PASS: --no-accel leaves the accelerator layer empty" + +# ⚠️ AND THE VARIABLE, WHICH THIS SECTION CLAIMED AND DID NOT MEASURE. +# `accel_str` prints `(none)` for an empty set so an ABI tag reads as a +# sentence, and handing that spelling on made `MCPP_ACCEL=(none)` reach every +# build program that asked for no accelerator — including every project that +# never mentioned one. The manual promised an empty string; a rule package +# testing "is there an accelerator" got a yes and a backend named `(none)`. +grep -q "accel=\[\]" none.log || { + cat none.log; echo "FAIL: MCPP_ACCEL is not empty under --no-accel"; exit 1; } +echo "PASS: --no-accel leaves MCPP_ACCEL empty, not the display spelling" + +echo "PASS: probe channel and accel reach the build program" diff --git a/tests/e2e/606_constrained_source_globs_narrow_to_the_build.sh b/tests/e2e/606_constrained_source_globs_narrow_to_the_build.sh new file mode 100755 index 000000000..0aa9e3825 --- /dev/null +++ b/tests/e2e/606_constrained_source_globs_narrow_to_the_build.sh @@ -0,0 +1,96 @@ +#!/usr/bin/env bash +# requires: gcc +# A `[build] sources` entry may carry the accel it is for, and the build +# narrows to it. +# +# sources = ["src/**/*.cpp", { glob = "src/dev/**/*", accel = "widget9+{w1}" }] +# +# Three outcomes, each measured here with its control: +# - the build covers the constraint: the glob's C++ files compile, and its +# device-kind files reach the build program as MCPP_DEVICE_SOURCES; +# - `--no-accel`: the glob is left out -- its C++ file is not compiled and +# the device list is empty. This is the CPU-only variant of one project; +# - the build targets something else: refused, naming the glob and both +# sides (`accel-mismatch`); +# - a constrained glob matching no file: refused, naming the glob. +# +# `widget` is not a backend anything knows. The engine reads a shape. +set -e + +TMP=$(mktemp -d) +trap "rm -rf $TMP" EXIT +cd "$TMP" + +"$MCPP" new narrow > /dev/null; cd narrow +rm -f src/*.cppm; mkdir -p src/dev +cat > src/main.cpp <<'EOF' +int main() { return 0; } +EOF +# A C++ file under the constrained glob: whether it is compiled is visible in +# the verbose command lines. A device-kind file beside it: the engine has no +# rule for it, and hands it to the build program. +printf 'int widget_helper() { return 42; }\n' > src/dev/helper.cpp +printf '__global__ void k() {}\n' > src/dev/kernel.cu + +write_manifest() { # $1 = accel of the build, $2 = accel of the glob, $3 = glob + cat > mcpp.toml < build.mcpp <<'EOF' +import mcpp; +int main() { + const char* d = mcpp::device_sources(); + mcpp::warning((*d ? d : "(no device sources)")); + return 0; +} +EOF +} + +# ── One: the build covers the constraint ───────────────────────────────── +write_manifest "widget9+{w1,w2}" "widget9+{w1}" "src/dev/**/*" +"$MCPP" build -v > covered.log 2>&1 || { cat covered.log; echo "FAIL: a covered constraint failed the build"; exit 1; } +grep -q "helper.cpp" covered.log || { cat covered.log; echo "FAIL: the constrained glob's C++ file was not compiled"; exit 1; } +grep -q "src/dev/kernel.cu" covered.log || { cat covered.log; echo "FAIL: the device source did not reach the build program"; exit 1; } +echo "PASS: a covered constraint compiles its C++ and hands its device sources to the build program" + +# ── Two: --no-accel leaves the glob out ─────────────────────────────────── +"$MCPP" build -v --no-accel > none.log 2>&1 || { cat none.log; echo "FAIL: --no-accel failed the build"; exit 1; } +if grep -q "helper.cpp" none.log; then + cat none.log; echo "FAIL: the constrained glob was compiled under --no-accel"; exit 1 +fi +grep -q "(no device sources)" none.log || { cat none.log; echo "FAIL: device sources were listed under --no-accel"; exit 1; } +echo "PASS: --no-accel leaves the constrained glob out, C++ and device files alike" + +# ── Three: the build targets something the constraint is not within ────── +write_manifest "widget9+{w2}" "widget9+{w1}" "src/dev/**/*" +if "$MCPP" build > mismatch.log 2>&1; then + cat mismatch.log; echo "FAIL: a constraint outside the build's accel was accepted"; exit 1 +fi +grep -q "src/dev/\*\*/\*" mismatch.log || { cat mismatch.log; echo "FAIL: refusal does not name the glob"; exit 1; } +grep -q "widget9+{w1}" mismatch.log || { cat mismatch.log; echo "FAIL: refusal does not name the constraint"; exit 1; } +grep -q "widget9+{w2}" mismatch.log || { cat mismatch.log; echo "FAIL: refusal does not name what the build targets"; exit 1; } +reason="$("$MCPP" why toolchain --format json 2>/dev/null | jq -r '.data.reason // "-"' | tr -d '\r')" +[[ "$reason" == "accel-mismatch" ]] || { echo "FAIL: reason is '$reason', expected accel-mismatch"; exit 1; } +echo "PASS: a constraint outside the build is refused naming the glob and both sides" + +# ── Four: a constrained glob that matches nothing ───────────────────────── +write_manifest "widget9+{w1,w2}" "widget9+{w1}" "src/nowhere/**/*" +if "$MCPP" build > empty.log 2>&1; then + cat empty.log; echo "FAIL: a constrained glob matching nothing was accepted"; exit 1 +fi +grep -q "src/nowhere" empty.log || { cat empty.log; echo "FAIL: refusal does not name the empty glob"; exit 1; } +echo "PASS: a constrained glob matching nothing is refused" + +echo "PASS: constrained source globs narrow to the build" diff --git a/tests/e2e/607_chained_actions_form_the_device_link.sh b/tests/e2e/607_chained_actions_form_the_device_link.sh new file mode 100755 index 000000000..cb304f7db --- /dev/null +++ b/tests/e2e/607_chained_actions_form_the_device_link.sh @@ -0,0 +1,85 @@ +#!/usr/bin/env bash +# requires: gcc +# One action may consume what another produced, and the engine only orders and +# fingerprints them. That is the whole engine-side content of a device link +# (multi-device design C-1/C-2): a rule package emits N `artifact` actions +# whose outputs stay out of the link, and one `object` action that reads them +# and produces the object that does join it. Nothing here names a device: +# the "device link" is `cat`, the "device compiler" is the toolchain's own C +# compiler, and the property under test is the graph, not the vendor. +# +# step 1 (artifact): src/parts/a.inc + src/parts/b.inc -> out/joined.c +# step 2 (object): out/joined.c -> out/joined.o (linked) +# +# Also measured: editing an input of step 1 rebuilds step 2 and relinks, which +# is the "changing the device link re-prepares" half of criterion C9. +set -e + +TMP=$(mktemp -d) +trap "rm -rf $TMP" EXIT +cd "$TMP" + +"$MCPP" new chain > /dev/null; cd chain +rm -f src/*.cppm; mkdir -p src/parts +cat > src/main.cpp <<'EOF2' +extern "C" int joined_value(); +int main() { return joined_value() == 42 ? 0 : 1; } +EOF2 +printf 'extern "C" int joined_value() { return\n' > src/parts/a.inc +printf '42; }\n' > src/parts/b.inc +cat > mcpp.toml <<'EOF2' +[package] +name = "chain" +version = "0.1.0" +[language] +standard = "c++23" +[targets.chain] +kind = "bin" +main = "src/main.cpp" +EOF2 +cat > build.mcpp <<'EOF2' +import std; +import mcpp; +int main() { + mcpp::rerun_if_changed("src/parts/a.inc"); + mcpp::rerun_if_changed("src/parts/b.inc"); + const std::string root = mcpp::manifest_dir(), out = mcpp::out_dir(); + const std::string joined = out + "/joined.cpp", obj = out + "/joined.o"; + { // Step 1: an ARTIFACT. Its output is data as far as the link is concerned. + mcpp::action a; + a.id = "join"; a.role = "artifact"; a.description = "join parts"; + a.arg("sh"); a.arg("-c"); + a.arg(("cat '" + root + "/src/parts/a.inc' '" + root + "/src/parts/b.inc' > '" + joined + "'").c_str()); + a.input((root + "/src/parts/a.inc").c_str()); + a.input((root + "/src/parts/b.inc").c_str()); + a.output(joined.c_str()); + a.submit(); + } + { // Step 2: an OBJECT that consumes step 1's output. The engine sees a + // path; ninja orders the two by it. + mcpp::action a; + a.id = "compile-joined"; a.role = "object"; a.description = "compile joined"; + a.arg((std::string(mcpp::toolchain_dir()) + "/bin/g++").c_str()); + a.arg("-c"); a.arg(joined.c_str()); a.arg("-o"); a.arg(obj.c_str()); + a.input(joined.c_str()); + a.output(obj.c_str()); + a.submit(); + } + return 0; +} +EOF2 + +"$MCPP" build > build1.log 2>&1 || { cat build1.log; echo "FAIL: the chained build failed"; exit 1; } +"$MCPP" run > run1.log 2>&1 || { cat run1.log; echo "FAIL: the linked object did not carry the joined value"; exit 1; } +echo "PASS: an object action consumed an artifact action's output and joined the link" + +# The second half of C9: a change at the head of the chain propagates. The +# joined value becomes 43, main returns 1, and only a relink would notice. +printf '43; }\n' > src/parts/b.inc +"$MCPP" build > build2.log 2>&1 || { cat build2.log; echo "FAIL: rebuild after editing an input failed"; exit 1; } +if "$MCPP" run > run2.log 2>&1; then + cat build2.log run2.log; echo "FAIL: the edit at the head of the chain did not reach the link"; exit 1 +fi +echo "PASS: editing an input of the first action rebuilt the second and relinked" + +echo "PASS: chained actions form the device link" diff --git a/tests/e2e/608_object_actions_reach_a_static_library.sh b/tests/e2e/608_object_actions_reach_a_static_library.sh new file mode 100755 index 000000000..ff1483bb7 --- /dev/null +++ b/tests/e2e/608_object_actions_reach_a_static_library.sh @@ -0,0 +1,84 @@ +#!/usr/bin/env bash +# requires: gcc +# An `object`-role action's outputs join a STATIC LIBRARY, not only an +# executable. That is the engine half of the multi-device design's C-6: a +# package whose device code is its point declares `kind = "lib"`, and until +# this the actions it emitted were dropped with a warning and the archive came +# out with none of them in it. +# +# Nothing here names a device. The "device compiler" is the toolchain's own C +# compiler and the property under test is which link units an action attaches +# to, which is not a vendor question. +set -e + +TMP=$(mktemp -d) +trap "rm -rf $TMP" EXIT +cd "$TMP" + +"$MCPP" new archived > /dev/null; cd archived +rm -f src/*.cppm src/main.cpp +mkdir -p src +cat > src/host.c <<'EOF2' +extern int from_action(void); +int host_value(void) { return from_action(); } +EOF2 +cat > mcpp.toml <<'EOF2' +[package] +name = "archived" +version = "0.1.0" +[language] +standard = "c++23" +[build] +sources = ["src/*.c"] +[targets.archived] +kind = "lib" +EOF2 +cat > build.mcpp <<'EOF2' +import std; +import mcpp; +int main() { + const std::string out = std::string(mcpp::out_dir()); + const std::string src = out + "/piece.c"; + { std::ofstream f(src, std::ios::trunc); + f << "int from_action(void) { return 7; }\n"; } + const std::string obj = out + "/piece.o"; + mcpp::action a; + a.id = "piece"; a.role = "object"; a.description = "compile the piece"; + a.arg((std::string(mcpp::toolchain_dir()) + "/bin/gcc").c_str()); + a.arg("-c"); a.arg(src.c_str()); a.arg("-o"); a.arg(obj.c_str()); + a.input(src.c_str()); + a.output(obj.c_str()); + a.submit(); + return 0; +} +EOF2 + +# ── One: the action runs and its object is IN the archive ──────────────── +"$MCPP" build > build.log 2>&1 || { cat build.log; echo "FAIL: the build failed"; exit 1; } + +grep -q "produces no target" build.log && { + cat build.log + echo "FAIL: the action was dropped for want of a linked image" + exit 1 +} + +lib=$(find target -name 'libarchived.a' | head -1) +[ -n "$lib" ] || { echo "FAIL: no static library was produced"; exit 1; } + +# ⭐ THE MEMBER LIST IS THE ASSERTION, not the exit status: an `ar` that was +# handed nothing still writes a well-formed archive and reports success. +members=$(ar t "$lib" 2>/dev/null | tr '\n' ' ') +case "$members" in + *piece.o*) ;; + *) echo "FAIL: piece.o is not a member of the archive (members: $members)"; exit 1 ;; +esac +echo "PASS: an object action's output is archived into a static library" + +# ── Two: the symbol is really there ───────────────────────────────────── +if command -v nm > /dev/null 2>&1; then + nm "$lib" 2>/dev/null | grep -q "from_action" || { + echo "FAIL: the archive has the member but not its symbol"; exit 1; } + echo "PASS: the archived member carries its symbol" +fi + +echo "PASS: object actions reach a static library" diff --git a/tests/unit/test_build_directives.cpp b/tests/unit/test_build_directives.cpp index f84e53d9c..a88c1a065 100644 --- a/tests/unit/test_build_directives.cpp +++ b/tests/unit/test_build_directives.cpp @@ -206,7 +206,9 @@ TEST(BuildDirectives, SerializeDeserializeRoundTrip) { "mcpp:generated=src/gen.cpp\n" "mcpp:source=vendor/pick.cpp\n" "mcpp:include-dir=inc\n" - "mcpp:include-dir-after=after\n"); + "mcpp:include-dir-after=after\n" + "mcpp:fact=widget.driver=1.2\n" + "mcpp:floor=widget.driver >= 1.0\n"); std::ostringstream os; dirs::serialize(os, d); @@ -231,6 +233,45 @@ TEST(BuildDirectives, SerializeDeserializeRoundTrip) { } } +// ── The probe channel ────────────────────────────────────────────────────── +// +// A rule package measures, the engine compares. What arrives on the wire is +// folded into the SAME runtime declarations a manifest could have written, +// so the version-floor check in prepare reads one list and never learns +// which spelling a claim came in. + +TEST(BuildDirectives, FactsAndFloorsAreClaimsThatFoldIntoRuntimeDeclarations) { + auto d = parse("mcpp:fact=widget.driver=1.2\n" + "mcpp:floor=widget.driver >= 2.0\n"); + EXPECT_EQ(d.at(dirs::Slot::Facts), (std::vector{"widget.driver=1.2"})); + EXPECT_EQ(d.at(dirs::Slot::Floors), (std::vector{"widget.driver >= 2.0"})); + + mcpp::manifest::Manifest m; + dirs::apply(m, d); + ASSERT_EQ(m.runtimeConfig.provides.size(), 1u); + EXPECT_EQ(m.runtimeConfig.provides[0], "widget.driver=1.2"); + ASSERT_EQ(m.runtimeConfig.requirements.size(), 1u); + EXPECT_EQ(m.runtimeConfig.requirements[0].kind, "version-floor"); + EXPECT_EQ(m.runtimeConfig.requirements[0].value, "widget.driver >= 2.0"); + // A claim about the build, not about the link or the run: the phase says + // when it is decided. + EXPECT_EQ(m.runtimeConfig.requirements[0].phase, "build"); +} + +TEST(BuildDirectives, AClaimReachesNeitherCompileNorLink) { + auto d = parse("mcpp:fact=widget.driver=1.2\n" + "mcpp:floor=widget.driver >= 2.0\n"); + mcpp::manifest::Manifest m; + dirs::apply(m, d); + EXPECT_TRUE(m.buildConfig.cflags.empty()); + EXPECT_TRUE(m.buildConfig.cxxflags.empty()); + EXPECT_TRUE(m.buildConfig.ldflags.empty()); + EXPECT_TRUE(m.buildConfig.sources.empty()); + for (auto const& def : dirs::kTable) + if (def.slot == dirs::Slot::Facts || def.slot == dirs::Slot::Floors) + EXPECT_EQ(def.scope, dirs::Scope::Claim) << def.wire; +} + TEST(BuildDirectives, RerunSlotsAreNotPersistedAsDirectives) { auto d = parse("mcpp:rerun-if-changed=config.h\n" "mcpp:rerun-if-env-changed=USE_FAST\n"); diff --git a/tests/unit/test_core_vendor_probes.cpp b/tests/unit/test_core_vendor_probes.cpp new file mode 100644 index 000000000..833b62438 --- /dev/null +++ b/tests/unit/test_core_vendor_probes.cpp @@ -0,0 +1,96 @@ +#include + +import std; + +// The engine owns no vendor probes. +// +// The bound a device toolkit states for its host compiler, whether a device +// compiler can reach its own back-end, and which driver a machine has are +// facts about one vendor's tools. They were once read in `src/doctor.cppm` and +// `src/toolchain/devicehost.cppm`, and every one of them was spelled CUDA. +// Moving them into the rule package that drives the tools is what keeps the +// engine from growing four copies -- AMD, Intel and Apple each have their own +// bound and their own back-end reachability question. +// +// What the engine keeps is the shape: the accel grammar (backend, version, +// architecture set, floor), the probe channel (`mcpp:fact` / `mcpp:floor`) +// and the comparison. None of those names a tool. +// +// The words below are assembled from pieces so this file does not flag itself +// should it ever move under src/. + +namespace { + +std::string without_comments(std::string_view source) { + enum class State { Code, Line, Block, String, Character } state = State::Code; + std::string out; + out.reserve(source.size()); + for (std::size_t i = 0; i < source.size(); ++i) { + const char c = source[i]; + const char n = i + 1 < source.size() ? source[i + 1] : '\0'; + if (state == State::Code) { + if (c == '/' && n == '/') { state = State::Line; ++i; out += " "; } + else if (c == '/' && n == '*') { state = State::Block; ++i; out += " "; } + else { + out.push_back(c); + if (c == '"') state = State::String; + else if (c == '\'') state = State::Character; + } + } else if (state == State::Line) { + if (c == '\n') { state = State::Code; out.push_back(c); } + else out.push_back(' '); + } else if (state == State::Block) { + if (c == '*' && n == '/') { state = State::Code; ++i; out += " "; } + else out.push_back(c == '\n' ? '\n' : ' '); + } else { + out.push_back(c); + if (c == '\\' && i + 1 < source.size()) out.push_back(source[++i]); + else if ((state == State::String && c == '"') + || (state == State::Character && c == '\'')) state = State::Code; + } + } + return out; +} + +} // namespace + +TEST(CoreVendorProbes, TheEngineNamesNoVendorTool) { + auto repo = std::filesystem::path(__FILE__).parent_path().parent_path().parent_path(); + auto src = repo / "src"; + ASSERT_TRUE(std::filesystem::is_directory(src)) << src; + // Tools, probes and locations. Each one is something only a rule package + // has business invoking or reading. + const std::vector tools = { + "nv" "cc", "ci" "cc", "pt" "xas", "fat" "binary", "nvidia" "-smi", + "cuDriver" "GetVersion", "/usr/local/" "cuda", "host_config" ".h", + "nvcc" ".profile", "hip" "cc", "rocm" "info", "ic" "px", "sycl" "-ls", + "cudafe" "++", + }; + std::size_t files = 0; + for (auto it = std::filesystem::recursive_directory_iterator(src); + it != std::filesystem::recursive_directory_iterator{}; ++it) { + if (!it->is_regular_file()) continue; + auto ext = it->path().extension().string(); + if (ext != ".cpp" && ext != ".cppm") continue; + ++files; + std::ifstream input(it->path()); + std::string raw((std::istreambuf_iterator(input)), {}); + auto code = without_comments(raw); + std::ranges::transform(code, code.begin(), + [](unsigned char c) { return std::tolower(c); }); + std::size_t line = 0; + for (auto text : code | std::views::split('\n')) { + ++line; + std::string lineText(text.begin(), text.end()); + for (auto const& tool : tools) { + std::string lowered = tool; + std::ranges::transform(lowered, lowered.begin(), + [](unsigned char c) { return std::tolower(c); }); + EXPECT_EQ(lineText.find(lowered), std::string::npos) + << it->path() << ':' << line << " names a vendor tool: " << tool; + } + } + } + // The denominator: a scan that found no files would pass vacuously. + EXPECT_GT(files, 100u); +} diff --git a/tests/unit/test_devicehost.cpp b/tests/unit/test_devicehost.cpp deleted file mode 100644 index 6ae8680ec..000000000 --- a/tests/unit/test_devicehost.cpp +++ /dev/null @@ -1,149 +0,0 @@ -#include - -import std; -import mcpp.toolchain.devicehost; - -using mcpp::toolchain::parse_host_config; -using mcpp::toolchain::host_compiler_accepted; - -namespace { -// The two guards as CUDA 12.0 writes them, reduced to what the parse reads. -constexpr std::string_view kCuda120 = R"( -#if __GNUC__ > 12 -#error -- unsupported GNU version! gcc versions later than 12 are not supported! The nvcc flag '-allow-unsupported-compiler' can be used to override this version check -#endif /* __GNUC__ > 12 */ -#if defined(__clang__) -#error -- unsupported clang version! clang version must be less than 15 and greater than 3.2 . -#endif -)"; -} // namespace - -TEST(DeviceHost, ReadsBothGuardsOutOfTheVendorHeader) { - auto b = parse_host_config(kCuda120); - EXPECT_TRUE(b.known()); - EXPECT_EQ(b.gccMax, 12); - EXPECT_EQ(b.clangMax, 14); // "less than 15" is an exclusive bound -} - -TEST(DeviceHost, AcceptsWithinTheBoundAndRefusesAbove) { - auto b = parse_host_config(kCuda120); - EXPECT_TRUE (host_compiler_accepted(b, "gcc", 12)); - EXPECT_FALSE(host_compiler_accepted(b, "gcc", 13)); - EXPECT_FALSE(host_compiler_accepted(b, "gcc", 16)); // mcpp's own payload - EXPECT_TRUE (host_compiler_accepted(b, "clang", 14)); - EXPECT_FALSE(host_compiler_accepted(b, "clang", 18)); - EXPECT_TRUE (host_compiler_accepted(b, "llvm", 14)); // mcpp's family name -} - -TEST(DeviceHost, AnUnreadableHeaderMakesNoClaim) { - // A refusal invented from a file the parse did not understand would be - // worse than the failure it prevents: the user cannot act on it. - auto b = parse_host_config("nothing to see here"); - EXPECT_FALSE(b.known()); - EXPECT_TRUE(host_compiler_accepted(b, "gcc", 99)); - EXPECT_TRUE(host_compiler_accepted(b, "clang", 99)); -} - -TEST(DeviceHost, SilenceAboutOneFamilyIsNotARefusalOfIt) { - auto b = parse_host_config("#if __GNUC__ > 11\n#error nope\n#endif\n"); - EXPECT_EQ(b.gccMax, 11); - EXPECT_EQ(b.clangMax, 0); - EXPECT_FALSE(host_compiler_accepted(b, "gcc", 12)); - EXPECT_TRUE (host_compiler_accepted(b, "clang", 20)); -} - -TEST(DeviceHost, AnUnknownVersionMakesNoClaimEither) { - auto b = parse_host_config(kCuda120); - EXPECT_TRUE(host_compiler_accepted(b, "gcc", 0)); -} - -// ── The plan nvcc states, and the one thing that goes missing from it ────── -// -// Both fixtures are real `nvcc --dryrun` output, CUDA 12.0, with the -// temporary paths shortened. They differ in exactly one line: the working -// host states a PATH, and the sandbox -- whose /etc is replaced, so the -// `nvcc.profile` symlinked into it is gone -- states none. Every other line, -// including the stages nvcc will invoke, is identical. - -using mcpp::toolchain::parse_dryrun; - -namespace { - -constexpr std::string_view kPlanWithProfile = R"(#$ _NVVM_BRANCH_=nvvm -#$ _SPACE_= -#$ _HERE_=/usr/lib/nvidia-cuda-toolkit/bin -#$ _TARGET_SIZE_=64 -#$ NVVMIR_LIBRARY_DIR=/usr/lib/nvidia-cuda-toolkit/libdevice -#$ PATH=/usr/lib/nvidia-cuda-toolkit/bin:/usr/local/bin:/usr/bin:/bin -#$ LIBRARIES= -L/usr/lib/x86_64-linux-gnu/stubs -#$ gcc -D__CUDA_ARCH_LIST__=520 -E -x c++ -m64 "/tmp/X" -o "/tmp/X" -#$ cudafe++ --c++17 --gnu_version=130300 --m64 "/tmp/X" -#$ cicc --c++17 -arch compute_52 -m64 "/tmp/X" -o "/tmp/X" -#$ ptxas -arch=sm_52 -m64 "/tmp/X" -o "/tmp/X" -#$ fatbinary -64 --cicc-cmdline="-ftz=0 " "--image3=kind=elf,sm=52,file=/tmp/X" -#$ rm /tmp/X -)"; - -// The same run with the PATH assignment removed: what nvcc emits when it -// cannot read its own profile. -constexpr std::string_view kPlanWithoutProfile = R"(#$ _NVVM_BRANCH_=nvvm -#$ _SPACE_= -#$ _HERE_=/usr/lib/nvidia-cuda-toolkit/bin -#$ _TARGET_SIZE_=64 -#$ LIBRARIES= -L/usr/lib/x86_64-linux-gnu/stubs -#$ gcc -D__CUDA_ARCH_LIST__=520 -E -x c++ -m64 "/tmp/X" -o "/tmp/X" -#$ cudafe++ --c++17 --gnu_version=130300 --m64 "/tmp/X" -#$ cicc --c++17 -arch compute_52 -m64 "/tmp/X" -o "/tmp/X" -#$ ptxas -arch=sm_52 -m64 "/tmp/X" -o "/tmp/X" -#$ fatbinary -64 --cicc-cmdline="-ftz=0 " "--image3=kind=elf,sm=52,file=/tmp/X" -#$ rm /tmp/X -)"; - -} // namespace - -TEST(DeviceDryRun, CollectsTheStagesAndThePathNvccStates) { - auto plan = parse_dryrun(kPlanWithProfile); - EXPECT_EQ(plan.searchPath, - "/usr/lib/nvidia-cuda-toolkit/bin:/usr/local/bin:/usr/bin:/bin"); - EXPECT_EQ(plan.programs, - (std::vector{"gcc", "cudafe++", "cicc", "ptxas", - "fatbinary", "rm"})); -} - -TEST(DeviceDryRun, TheMissingProfileShowsUpAsAnAbsentPathAssignment) { - // This is the whole of the difference the check keys on. The stages are - // the same; only the path they will be resolved against is gone. - auto broken = parse_dryrun(kPlanWithoutProfile); - auto intact = parse_dryrun(kPlanWithProfile); - EXPECT_TRUE(broken.searchPath.empty()); - EXPECT_EQ(broken.programs, intact.programs); -} - -TEST(DeviceDryRun, AssignmentsAreNotMistakenForStages) { - auto plan = parse_dryrun(kPlanWithProfile); - // `LIBRARIES= -L...` and `_SPACE_= ` both parse as assignments, and - // neither names a program. Reading either as a stage would report a - // missing tool that nvcc never intended to run. - for (auto const& p : plan.programs) { - EXPECT_EQ(p.find('='), std::string::npos); - EXPECT_NE(p, "LIBRARIES"); - EXPECT_NE(p, "_SPACE_"); - } -} - -TEST(DeviceDryRun, StagesNamedByPathAreLeftAlone) { - // A stage nvcc spells out resolves without the search path, so it is not - // a candidate for "cannot be found on PATH". - auto plan = parse_dryrun("#$ PATH=/bin\n" - "#$ /opt/cuda/bin/cicc --c++17\n" - "#$ ptxas -arch=sm_52\n"); - EXPECT_EQ(plan.programs, (std::vector{"ptxas"})); -} - -TEST(DeviceDryRun, TextThatIsNotAPlanYieldsNoStages) { - // A spawn that failed because there is no nvcc lands here. No stages - // means no finding: the probe reached no answer and invents none. - auto plan = parse_dryrun("mcpp: failed to spawn 'nvcc': No such file\n"); - EXPECT_TRUE(plan.programs.empty()); - EXPECT_TRUE(plan.searchPath.empty()); -} diff --git a/tests/unit/test_manifest.cpp b/tests/unit/test_manifest.cpp index 965a1b0ef..6554670a9 100644 --- a/tests/unit/test_manifest.cpp +++ b/tests/unit/test_manifest.cpp @@ -4862,6 +4862,93 @@ hardware = {} // the binding-time check is the one that acts on it and a warning must not // silently change what was declared. +// ── [build] sources: constrained entries ─────────────────────────────────── +// +// A table entry carries a glob and the accel it is for. The glob joins the +// plain list (every existing reader keeps working); the constraint is what +// prepare narrows on. + +TEST(ManifestSourceConstraints, ATableEntryJoinsTheListAndCarriesItsConstraint) { + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "infer" +version = "0.1.0" +[build] +sources = ["src/**/*.cppm", { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }] +)"); + ASSERT_TRUE(m.has_value()) << m.error().format(); + ASSERT_EQ(m->buildConfig.sources.size(), 2u); + EXPECT_EQ(m->buildConfig.sources[0], "src/**/*.cppm"); + EXPECT_EQ(m->buildConfig.sources[1], "src/kernels/**/*.cu"); + EXPECT_TRUE(m->buildConfig.sourcesDeclared); + ASSERT_EQ(m->buildConfig.sourceConstraints.size(), 1u); + EXPECT_EQ(m->buildConfig.sourceConstraints[0].glob, "src/kernels/**/*.cu"); + EXPECT_EQ(m->buildConfig.sourceConstraints[0].accel, "cuda12.9+{sm_89}"); +} + +TEST(ManifestSourceConstraints, ATableEntryWithoutAccelIsJustAGlob) { + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "infer" +version = "0.1.0" +[build] +sources = [{ glob = "src/**/*.cpp" }] +)"); + ASSERT_TRUE(m.has_value()) << m.error().format(); + ASSERT_EQ(m->buildConfig.sources.size(), 1u); + EXPECT_TRUE(m->buildConfig.sourceConstraints.empty()); +} + +TEST(ManifestSourceConstraints, AnEntryWithoutGlobIsRefused) { + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "infer" +version = "0.1.0" +[build] +sources = [{ accel = "cuda12.9+{sm_89}" }] +)"); + ASSERT_FALSE(m.has_value()); + EXPECT_NE(m.error().format().find("glob"), std::string::npos) << m.error().format(); +} + +TEST(ManifestSourceConstraints, AnUnknownKeyIsRefusedRatherThanSkipped) { + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "infer" +version = "0.1.0" +[build] +sources = [{ glob = "src/k/*.cu", acel = "cuda12.9+{sm_89}" }] +)"); + ASSERT_FALSE(m.has_value()); + EXPECT_NE(m.error().format().find("acel"), std::string::npos) << m.error().format(); +} + +TEST(ManifestSourceConstraints, AnExclusionCannotBeConstrained) { + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "infer" +version = "0.1.0" +[build] +sources = [{ glob = "!src/k/*.cu", accel = "cuda12.9+{sm_89}" }] +)"); + ASSERT_FALSE(m.has_value()); +} + +TEST(ManifestSourceConstraints, AConstrainedEntryUnderTheConditionalAxisIsRefused) { + // read_list keeps strings only; a table that vanished there would be a + // device glob nobody ever narrowed, so it is refused with a pointer to + // where it belongs. + auto m = mcpp::manifest::parse_string(R"( +[package] +name = "infer" +version = "0.1.0" +[target.'cfg(linux)'.build] +sources = [{ glob = "src/k/*.cu", accel = "cuda12.9+{sm_89}" }] +)"); + ASSERT_FALSE(m.has_value()); + EXPECT_NE(m.error().format().find("[build].sources"), std::string::npos) << m.error().format(); +} + TEST(ManifestExclusive, CarriesTheListAndWarnsAboutNothing) { auto m = mcpp::manifest::parse_string(R"( [package] diff --git a/tests/unit/test_version_floor.cpp b/tests/unit/test_version_floor.cpp new file mode 100644 index 000000000..f409d2abd --- /dev/null +++ b/tests/unit/test_version_floor.cpp @@ -0,0 +1,73 @@ +#include + +import std; +import mcpp.build.version_floor; + +using mcpp::build::parse_version_floor; +using mcpp::build::parse_version_fact; +using mcpp::build::version_at_least; + +// ── The two spellings ────────────────────────────────────────────────────── + +TEST(VersionFloor, ReadsAFloor) { + auto f = parse_version_floor("cuda.driver >= 12.0"); + ASSERT_TRUE(f.valid()); + EXPECT_EQ(f.name, "cuda.driver"); + EXPECT_EQ(f.version, "12.0"); +} + +TEST(VersionFloor, ReadsAFact) { + auto f = parse_version_fact("cuda.driver=12.4"); + ASSERT_TRUE(f.valid()); + EXPECT_EQ(f.name, "cuda.driver"); + EXPECT_EQ(f.version, "12.4"); +} + +TEST(VersionFloor, AFloorIsNotReadAsAFact) { + // Both spellings live in string lists, and `>=` contains `=`. Reading one + // as the other would turn "needs at least 12.0" into "this machine has + // 12.0" — a requirement silently becoming its own satisfaction. + EXPECT_FALSE(parse_version_fact("cuda.driver >= 12.0").valid()); +} + +TEST(VersionFloor, AnythingElseIsNotAFloor) { + EXPECT_FALSE(parse_version_floor("cuda.driver").valid()); + EXPECT_FALSE(parse_version_floor("cuda.driver >= ").valid()); + EXPECT_FALSE(parse_version_floor(" >= 12.0").valid()); + // Not a version: refused rather than read as `12`. + EXPECT_FALSE(parse_version_floor("cuda.driver >= 12.").valid()); + EXPECT_FALSE(parse_version_floor("cuda.driver >= twelve").valid()); +} + +// ── Comparison ───────────────────────────────────────────────────────────── + +TEST(VersionFloor, ComparesComponentwise) { + EXPECT_EQ(version_at_least("12.4", "12.0"), std::optional{true}); + EXPECT_EQ(version_at_least("12.0", "12.4"), std::optional{false}); + EXPECT_EQ(version_at_least("13.0", "12.9"), std::optional{true}); + // Not lexicographic: 12.10 is above 12.9. + EXPECT_EQ(version_at_least("12.10", "12.9"), std::optional{true}); +} + +TEST(VersionFloor, AMissingComponentIsZero) { + EXPECT_EQ(version_at_least("12", "12.0"), std::optional{true}); + EXPECT_EQ(version_at_least("12.0", "12"), std::optional{true}); + EXPECT_EQ(version_at_least("12", "12.1"), std::optional{false}); +} + +TEST(VersionFloor, EitherSideUnreadableIsNoClaim) { + // The whole reason this mechanism exists is that a wrong answer is worse + // than no answer. A version that cannot be read yields neither. + EXPECT_FALSE(version_at_least("unknown", "12.0").has_value()); + EXPECT_FALSE(version_at_least("12.0", "").has_value()); +} + +TEST(VersionFloor, NoVendorVocabularyIsRequired) { + // The name is data. A backend this engine has never heard of compares the + // same way, which is what keeps the comparison out of the vendors' reach. + auto f = parse_version_floor("some.future.thing >= 4.2.1"); + ASSERT_TRUE(f.valid()); + EXPECT_EQ(f.name, "some.future.thing"); + EXPECT_EQ(version_at_least("4.3", f.version), std::optional{true}); + EXPECT_EQ(version_at_least("4.2", f.version), std::optional{false}); +}