diff --git a/.agents/docs/2026-09-05-accelerator-support-design.md b/.agents/docs/2026-09-05-accelerator-support-design.md new file mode 100644 index 000000000..9d0b69861 --- /dev/null +++ b/.agents/docs/2026-09-05-accelerator-support-design.md @@ -0,0 +1,1215 @@ +# 加速器支持:完整设计方案 + +2026-09-05 · **v5,完整方案**,未实现,供 review + +前置:[生态调研报告](2026-09-04-ai-accelerator-toolchain-ecosystem-survey.md) · +取代 [`2026-09-05-accelerator-support-scenarios-and-design.md`](2026-09-05-accelerator-support-scenarios-and-design.md)(v1/v2) + +对现状的陈述一律读 `origin/main`(d18cd8b,2026-09-04)。 +**[已有]** = 当前代码里就有;**[提议]** = 本文要加的。 + +--- + +## 0. 决定记录 + +| # | 决定 | 落在 | +|---|---|---| +| 1 | `accel` **不进 tag 字符串**,作为并列结构化字段,仍由 `tag_check` 一并比较 | §7.1 | +| 2 | `kind = "device"` + `backend`,不用 `language = "cuda"` | §3.1 | +| 3 | 逐 glob 架构收窄**第一版就做** | §3.2 | +| 4 | capability 同名符号边界**先修** | §7.4 | +| 5 | cfg 用**组合表达**,不引入第二个谓词 | §3.4 | +| 6 | 第一版 device target **不产静态库** | §6.3 | +| 7 | 岛的 C++ 运行时耦合**必须核验**,复用 `mcpp.runtime.elf` | §7.3 | +| 8 | 形态 B 的工具链声明**沿用 mcpp 既有的 `family@version` 写法** | §4 R8 | +| 9 | 小众后端(需注册的国产三家、编模型的那一类)**本轮不支持** | §1.2 | +| 10 | 方案必须含**生态闭环**:索引侧的包、真实验证、真实可用 | §8 | +| 11 | `archs` 缺省 ⇒ **报错并列出该 backend 的已知取值**,不猜 | §3.2 | +| 12 | `backend` 是**请求**(第一趟可见),绑定在第二趟,缝上有显式 satisfied 记录 | §3.5 | +| 13 | 形态 B 的产物核验 = **引擎读节名,节名表由后端描述符提供** | §7.5 | +| 14 | Metal 一类产物走**已有的 `deploy_files` 通道**,不新增产物类别 | §6.4 | +| 15 | 形态 B 的字段**统一叫 `archs`**,与形态 A 一致 | §4 R8 | +| 16 | **不加**「已构建未运行验证」标注 —— 交叉编译产物本来就没在目标上跑过 | §9 RK-4 | + +### 0.1 v4 相对 v3 的改动 + +- **§4 R8**:`[toolchain] compiler = "icpx"` 是我发明的写法,**错了**。mcpp 既有形态是 + `[toolchain] default = "gcc@16.1.0"` / `linux = "…"` / `macos = "…"`,值为 + `family@version`。⇒ SYCL 应当是 `default = "oneapi@2026.1"` 或 `"adaptivecpp@25.02"` + —— **oneAPI 与 AdaptiveCpp 是工具链家族,和 gcc/llvm/msvc 平级。** +- **§1.2**:范围收窄,小众后端明确不做。 +- **新增 §8 生态闭环**、**新增 §9 风险**。自我 review(§10)按新内容重写。 + +### 0.2 v5 相对 v4 的改动 + +- **RK-2 由代码给出了答案,风险等级下调**,但暴露出一个不同的残余风险(§9 RK-2)。 +- **新增 §8.5「没有 GPU 能验到哪一步」** —— 部分后端的 kernel **可以在 CPU 上真跑**, + 这改变了 AdaptiveCpp 优先的理由,也缩小了 RK-1。 +- **新增 §11 源码分发与二进制分发**。 +- v4 的六条待决全部拍板(决定 11–16),`§12` 只剩真正未决的。 + +--- + +## 1. 范围 + +### 1.1 本轮覆盖 + +| 类别 | 形态 | 代表 | +|---|---|---| +| CUDA 系单源方言 | A(岛) | CUDA、HIP/ROCm | +| 单源整目标 | B | SYCL(oneAPI DPC++ / AdaptiveCpp)、OpenMP offload | +| 着色/计算语言 | A(产资源非目标码) | Metal、Vulkan/SPIR-V | +| 纯绑定层 | 都不是 | 厂商 runtime 的 shim,只是普通库 | + +### 1.2 本轮明确不做(决定 9) + +- **需注册或需联系厂商才能拿到工具链的**:天数智芯 CoreX、壁仞 BRCC、燧原 TopsCC。 + 包管理器无法供给的东西,做了也是半截。 +- **「编模型不编 C++」的那一类**:Google TPU、AWS Neuron 主线、Qualcomm QNN、 + ARM Ethos-U、Rockchip RKNN。报告 §5.4:那里根本没有 C++ 编译步骤。 +- **昇腾 Ascend C、寒武纪 BANG、摩尔线程 MUSA**:形态上属于 A,框架天然容得下 + (§3.1 给了 manifest 长什么样),但**本轮不做规则包与载荷**,留给生态或后续。 + +⭐ 但**框架不为它们开特例**:§3.1 会证明同一组原语能表达它们。 +「不做」是排期决定,不是设计缺口。说清楚边界本身是产品的一部分。 + +--- + +## 2. 两种形态 + +### 2.1 形态 A —— 岛 + +设备代码在独立编译单元里,由另一个编译器编译,产物汇进宿主构建。 + +``` +kernels.cu ──nvcc──→ kernels.o(内嵌设备镜像) ──┐ + ├─→ 宿主链接 +main.cppm ──g++───→ main.o ────────────────────┘ +``` + +⚠️ **岛的产物不一定是目标文件。** Metal 是 +`xcrun metal` → `.air`、`metal-ar` → `.metalar`、`metallib` → `.metallib`, +最终是**运行期资源**;Vulkan 的 `.spv` 同理。 + +⇒ 岛的产物角色必须可声明。**[已有]** `BuildAction::Role` 恰好有 `Object`(进链接) +与 `Artifact`(不进链接)两档,规则包声明用哪个。**框架零新增。** + +### 2.2 形态 B —— 整目标 + +设备代码就在普通 `.cpp` 里,由一个能 offload 的编译器整体编译。**没有岛。** + +``` +main.cpp ──icpx -fsycl -fsycl-targets=…──→ main.o(含 SPIR-V/PTX/amdgcn 镜像) +``` + +形态 B 里「加速器支持」= **换一个工具链家族 + 一组 target flag**,不是加一个 target。 +用到 P3 + P4 + 工具链选择,**不用 P1**。 + +### 2.3 ⚠️ 形态 B 与 modules 的关系是**已知的未知** + +形态 A 有确定答案:nvcc 不支持 modules,岛不进 module graph(报告 §6)。 + +形态 B 没有。`icpx` 是 clang 派生、支持 `-std=c++20`,但 +**「SYCL kernel 所在的 TU 能否是 module interface」本次调研找不到任何肯定或否定的文档** +(报告 §8 第 1 条)。AdaptiveCpp 同理。 + +⇒ 不假设可行,也不假设不可行。第一版按「该目标退回 header 模式」处理, +并把实测列为阶段 0 的探针(§11)。**这是本设计唯一一处已知的未知,单独标出而不藏在假设里。** + +--- + +## 3. 主体框架:五个原语 + +| # | 原语 | 是什么 | 形态 | 状态 | +|---|---|---|---|---| +| **P1** | device target | 不进 module graph 的编译单元集合 | A | **[提议]** | +| **P2** | 接缝模块 | 普通模块;收窄可见性/类型,**且是后端替换的落点** | A | **[已有]** | +| **P3** | `accel` 字段 | 产物身份:含哪些 (后端, 架构集合) | A+B | **[提议]** | +| **P4** | `accelerator` layer key | cfg 一档:解析后的后端身份 | A+B | **[提议]** | +| **P5** | 规则包 | 厂商 flag 拼法的持有者 | A | **[已有]** | + +零改动复用:`BuildAction::Role`、`tag_check` + `prebuilt.cppm` 的选择循环、 +capability 绑定、11 字段指纹、xlings 载荷、`mcpp.runtime.elf`、`CommandDialect`。 + +### 3.1 P1 —— device target,且不为任何厂商开特例(决定 2) + +```toml +[targets.kernels] +kind = "device" +backend = "cuda" # 可省略:只有一个 device 规则包时自动绑 +sources = ["src/kernels/*.cu"] +archs = ["sm_80", "sm_90f"] +``` + +**本轮不做、但框架已经容得下的**(证明同一组原语够用): + +```toml +# 华为昇腾:.asc + --npu-arch=dav-2201 +[targets.ops] kind="device"; backend="ascend"; sources=["src/*.asc"]; archs=["dav-2201"] + +# 寒武纪:.mlu + 可重复的 --bang-mlu-arch +[targets.ops] kind="device"; backend="bang"; sources=["src/*.mlu"]; archs=["mtp_372","mtp_290"] + +# Metal:产物走 Role::Artifact,由规则包声明 +[targets.shaders] kind="device"; backend="metal"; sources=["src/*.metal"] +``` + +⭐ `kind = "device"` 说的是**图上的角色**,不说语言。这一点在 Intel Gaudi 上兑现: +TPC-C 是 **C99 派生而非 C++**,`kind = "device"` 对它照样成立 —— +若当初选 `language = "cuda"` 式命名,这里就要开特例。 + +### 3.2 架构收窄:一个键(决定 3) + +目标级声明全集,glob 级**收窄**;收窄即求交,不需要第二个键。 + +```toml +[targets.myops] +kind = "device" +archs = ["sm_80", "sm_90f", "sm_100f"] + +[targets.myops.sources] +"src/gemm_sm90/*.cu" = { archs = ["sm_90f"] } +"src/gemm_common/*.cu" = { archs = ["sm_80", "sm_90f"] } +``` + +规则:glob 级必须是目标级的**子集**;不是子集是配置错误,不是隐式扩展。 + +⚠️ **收窄成空集是错误,不是「不编」。** 空集意味着这批源码在当前请求的架构下没有 +有效目标;静默跳过会产出缺 kernel 的库 —— 报告里 `no kernel image` 的又一个来源。 + +理由(报告 §4.2):vLLM / ONNX Runtime / TensorRT-LLM 三家独立且趋同地做了逐 kernel +家族的收窄,因为无条件为每个 SM 实例化每个 CUTLASS 模板,编译太慢、二进制太胖。 + +#### `archs` 缺省 ⇒ 报错,不猜(决定 11) + +``` +error: [targets.kernels] backend = "cuda" requires `archs`. + mcpp does not guess: the set you compile for is not the set this machine has. + known values for cuda : sm_80 sm_86 sm_89 sm_90 sm_90a sm_100 sm_100f sm_120 + family targets (preferred, one artifact covers a generation): + sm_90f sm_100f + or, for local iteration only: archs = ["native"] +``` + +三条理由,都指向同一个方向: + +1. ⭐ **mcpp 已有的先例是「拒绝并告诉你该写什么」**,不是猜。 + `[toolchain] … = "system"` 被明确拒绝,错误里直接给出该写 `linux = "gcc@16.1.0"`。 +2. 报告 §2.2 记的 CMake `native` 陷阱,根因就是**把「本机有什么」当成「要编成什么」**; + 猜一个默认就是把这个陷阱抄过来。 +3. 源码分发时(§11.1)一个错的默认值可能是一小时的编译。**猜错的代价不对称。** + +### 3.3 P2 —— 接缝模块 + +不是新原语,但是**架构单元**,因为它同时做三件事,而第三件决定它不能省: + +| 它做什么 | 没有它会怎样 | +|---|---| +| 可见性收窄:下游 `import`,不见设备头 | 下游被迫包含设备头 | +| 类型收窄:`const float*, size_t` → `std::span` | 裸指针泄漏到整个工程 | +| ⭐ **后端可替换**:换掉岛而接口一个字不变 | **`cfg` 无处安放,每个消费者都变成后端相关的** | + +边界处可以不写头文件(GMF 里直接 `extern "C"`),但 ⚠️ `extern "C"` 不做名字修饰 +⇒ **改了 `.cu` 的签名而声明不跟着改,链接照样成功,运行期读错内存**。 +⇒ 默认用头文件(签名只有一份拷贝);「由规则包生成声明」作为可选能力。 + +### 3.4 P4 —— `accelerator` 多值 layer key,组合表达(决定 5) + +**[已有]** `kCfgLayerKeys` 现有五项,由 `merge_layer_conditional_config` 的**第二趟** +在目标侧解析后求值,未知键有 `unknown_tokens()` 诊断。加 `accelerator` 是加一项。 + +⚠️ 但它必须是**多值** layer(现有五项都是单值),而语义只有一行: + +> **`accelerator = ""` ⟺ ` ∈` 集合。处处如此。** + +`any` / `all` / `not` 作为普通布尔组合子在其上组合,不改变操作数的含义。 +单后端构建的集合是 `{cuda}`,于是它对 `accelerator = "cuda"` 答真、 +对 `accelerator = "rocm"` 答假 —— 这正是 R2/R7 需要的; +多后端构建的集合是 `{cuda, rocm}`,两者都答真 —— 这正是 R5 需要的。 + +⚠️⚠️ **本节最初写的是另一套语义**(裸键表示集合相等、`any(...)` 才表示成员判定), +实施时发现它是错的:让组合子改变操作数的含义,会使 +`all(accelerator = "cuda", accelerator = "rocm")` 变成**不可满足**, +而不是「两个后端都启用」。改为处处成员判定之后,不但正确,而且 +**用户少学一条规则** —— R5 的 manifest 里也不再需要写 `any(...)`。 +详见 §14.1 第 1 条。实现里这个差异只落在一个函数(`Ctx::layer_matches`)上。 + +### 3.5 P5 —— 规则包,`backend` 走 capability + +**[已有]** 规则包机制 2026.8.5.1 落地,`examples/08-build-rules/` 在 main 上。 + +`backend = "cuda"` **不是包名**,是能力名: + +```toml +[package] +name = "rules-cuda" +provides = ["mcpp:device-rule=cuda"] +``` + +**[已有]** capability 绑定规则原样适用(恰好一个自动绑;多个未 pin 报错并列候选; +零个报错),**绝不静默猜测**。⭐ 用户声明**意图**而非包名。 + +#### `backend` 在第几趟解析(决定 12) + +`backend = "cuda"` 是**请求**,第一趟(triple-only)就能读到它这个字符串; +把它**绑定**到具体规则包要走 capability,而 capability 需要解析后的图 ⇒ **第二趟**。 + +⭐ 这不是新规矩,是 mcpp 已有的一条经验的直接套用 —— +memory `a-printed-value-promoted-to-a-compared-value` 记的: +**「请求侧与答案侧保留各自的拼写,在缝上放一个显式的 `*_request_satisfied`」**。 + +⇒ 具体化为三条: + +- `targets[].backend` 是**请求侧**的字符串,不做归一,第一趟可见; +- `mcpp:device-rule=` 的绑定结果是**答案侧**,第二趟产生; +- 缝上记一个显式的 `backend_request_satisfied`,**它有读者**(诊断与 `mcpp why`), + 否则就是 memory 里「答案已解析,却没有接到决定上」那一类。 + +这样也切开了 §10.1 提到的循环嫌疑:第一趟不需要知道哪个规则包会赢。 + +--- + +## 4. 七类使用者 × 两种形态 + +| # | 使用者 | 形态 | 原语 | 新原语? | +|---|---|---|---|---| +| R1 | 应用开发者:只想用 GPU 库 | — | P3 | 否 | +| R2 | 混合工程:C++ + kernel | A | P1+P2+P5 | P1 | +| R3 | 算子库作者:独立发布 | A | P1+P2+P3 | P1、P3 | +| R4 | 算子开发工作流 | A | P1(逐 glob 收窄) | 否 | +| R5 | 推理框架:一份产物多后端 | A×N | P1×N+P3(集合)+P4 | 否 | +| R6 | 底层运行时/绑定层 | — | P4 + capability | **零新增** | +| R7 | 换厂商/国产迁移 | A 或 B | P4 + capability + P2 | 否 | +| R8 | SYCL / OpenMP / stdpar | **B** | P3+P4+工具链家族 | 否 | + +### R1 —— 只想用一个 GPU 库(数量最大) + +```toml +[dependencies] +onnxruntime = "1.23.0" +``` + +**[已有]** 消费侧遍历 `runtimeConfig.artifacts`,逐个 `tag_check`,取第一个通过的; +全不通过报**分歧最少的那一个**: + +> Keep the CLOSEST refusal to show: the one that disagrees least is the one the +> user is most likely able to act on. + +**这就是 PEP 817 的选择算法,已经在跑。** + +``` +error: onnxruntime@1.23.0: no prebuilt artifact matches this build. + your build : x86_64-linux-gnu-gcc16-libstdcxx16-c++23 accel=cuda12.8+{sm86} + published : (no accel) ← CPU-only + accel=cuda12.8+{sm80,sm90f} ptx-floor=90 + closest differs on: + accel needs sm80|sm90f, or PTX floor ≤ 86; this build asks sm86 +note: the CPU-only variant is compatible — take it with --no-accel. +``` + +### R5 —— 一份产物含多个后端(ggml 形态) + +```toml +[features] +cuda = {} +rocm = {} # 不互斥:可同时开 + +[target.'cfg(any(accelerator = "cuda"))'.targets.k-cuda] +kind = "device"; backend = "cuda"; archs = ["sm_90f"] + +[target.'cfg(any(accelerator = "rocm"))'.targets.k-rocm] +kind = "device"; backend = "rocm"; archs = ["gfx942"] +``` + +⇒ `accel` 是**集合**,判据是「请求的 (后端, 架构) 在不在里面」; +单后端退化成集合大小为一,R1 语义不变。 + +### R8 —— SYCL / OpenMP(形态 B),沿用 mcpp 既有工具链写法(决定 8) + +**没有 device target。** 变的是工具链家族: + +```toml +# mcpp 既有写法:[toolchain] = "@" +[toolchain] +default = "adaptivecpp@25.02" # 或 "oneapi@2026.1" + +[build] +accel = { backend = "sycl", archs = ["generic"] } # [提议];字段名与形态 A 统一(决定 15) +``` + +```toml +# OpenMP offload:工具链仍是 gcc,但载荷必须带 offload +[toolchain] +default = "gcc@16.1.0+offload" # [提议] 载荷变体,见 §5.3 +[build] +accel = { backend = "openmp", archs = ["nvptx-none"] } +``` + +⭐ **oneAPI 与 AdaptiveCpp 是工具链家族,和 gcc/llvm/msvc 平级。** 于是 +`@version` 钉定、自动安装、`mcpp toolchain` 命令、载荷映射**全部复用现有机制**, +加一个家族 = 注册表加一行 + 载荷映射,与当初加 llvm 是同一件事。 + +⭐ 一个待验证的简化:icpx 与 AdaptiveCpp 都是 clang 派生,flag 用 GNU 风格 +(**[已有]** `CommandDialect` 的 `gnu` 实例可直接复用),BMI 也是 clang 的。 +⇒ **可能不需要给 `CompilerId` 加新值,只需要加工具链家族。** 列为阶段 0 的探针。 + +⚠️ OpenMP offload **不是加 flag,是载荷要重建** —— GCC 必须在构建时配置了 +`--enable-offload-targets=nvptx-none=…`。没有 offload 支持的 GCC 遇到 `-foffload=` +**只会安静地什么都不做** ⇒ 判据是载荷的构建配置,不是编译能否通过。 + +--- + +## 5. 工具链供给 + +### 5.0 DPC++ 与 NVHPC 怎么接进来(2026-09-05 追加,已查证) + +这一节回答「intel/llvm 与 NVIDIA 的编译器要不要也走 xim、像普通工具链一样」。 +答案不是同一个,而且查证之后**比原设计更简单**。 + +| 编译器 | 接法 | 理由(一手材料) | +|---|---|---| +| **intel/llvm DPC++** | **`Family::Llvm` 的一个载荷变体 + provider 能力位**,**不是新家族** | 它就是 clang:同一个 driver、同一套 GNU 风格 flag、同一种 BMI。用法是 `clang++ -fsycl` | +| **nvcc** | **device 规则包**(形态 A) | 不支持 C++20 modules(报告 §6) | +| **nvc++ (NVHPC)** | **device 规则包,不是工具链家族** | HPC SDK 用户指南自述主标准是 *"ISO/ANSI C++17 with GNU compatibility"*,未见 modules 支持 ⇒ **当不了 modules-first 工程的宿主编译器** | + +#### 两处对 v4 的更正 + +**其一,开源 DPC++ 不需要 Codeplay 插件。** §5.1 原写「NVIDIA/AMD 后端还要 Codeplay +插件」—— 那是 **Intel 二进制发行版**的情况。开源的 intel/llvm 用 +`buildbot/configure.py --cuda` / `--hip` 把后端**编进编译器本身**,可用的 +`-fsycl-targets` 直接包括 `nvptx64-nvidia-cuda`、`amdgcn-amd-amdhsa`。 +先决条件:CUDA ≥ 11.0(建议 11.6+)、ROCm ≥ 5.7。 +(来源:`intel/llvm` 的 `sycl/doc/GetStartedGuide.md`) + +**其二,DPC++ 另有一个 `native_cpu` 后端。** 与 AdaptiveCpp 的 omp flow 同类: +**SYCL kernel 可以在 CPU 上跑**,于是形态 B 有第二条可在无卡 CI 里端到端自证的路径。 + +#### ⭐⭐ 由此,形态 B 不需要新的 `CompilerId`,也不需要新的 `Family` + +`src/toolchain/registry.cppm` 的 `Family` 是封闭枚举(`Gcc`/`Llvm`/`Msvc`), +加一项要动约十处。**而 DPC++ 一处都不用动** —— 它是 `Llvm` 家族的另一个载荷, +差异落在两个已有的地方: + +1. **载荷映射**:registry 里 family→ximName 的映射扩成 family+变体→ximName; +2. **能力位**:`ProviderCapabilities`(`src/toolchain/provider.cppm`)加 + `has_sycl` 与可用的 `-fsycl-targets` 列表 —— 该模块的存在理由原文就是 + *"Previously these decisions were scattered as ad-hoc is_clang(tc) / is_gcc(tc) + checks. This module centralises them into a single query point."* + +⇒ 形态 B 的工程量从「加一个工具链家族」降到「加一个载荷变体 + 一个能力位」。 + +#### nvc++ 归形态 A,不归形态 B + +这条值得单独说,因为它反直觉:nvc++ 提供 `-stdpar`,看起来正是形态 B 的样子。 +但**形态 B 的前提是那个编译器能当整个工程的宿主编译器**,而 mcpp 的工程全都 +`import std`。一个主标准是 C++17、无 modules 的编译器编不了它们。 +⇒ nvc++ 只能作为**某些目标的设备编译器**出现,与 nvcc 同类,走规则包。 + +### 5.0b ⭐⭐ 宿主依赖:轴已经划好了,而且 xim 里已经有半边 + +设计原则(用户长期约束):**尽量走 xlings 生态,越少依赖 host 越好。** + +这条在本课题上不需要新规则 —— memory 里已有的轴划分逐字适用: + +> **工具链 = mcpp 的契约 ⇒ 拒绝 `system`;程序链接的库 = 程序自己的事 ⇒ warn 指回 mcpp-index** + +套到加速器上,轴是这样切的: + +| 组件 | 归属 | 理由 | +|---|---|---| +| `nvcc`、`ptxas`、`nvlink`、`fatbinary`、CRT 头 | **xim 载荷** | 它是工具链,是 mcpp 的契约。NVIDIA 的 redist 清单允许再分发 | +| `cudart`、`cublas`、`cudnn` 等运行库 | **xim 载荷**(mcpp-index 侧的 `compat.*` 消费) | 同样可再分发 | +| **`libcuda.so.1`(驱动 userspace)** | ⚠️ **不可分发,只能指向 host** | 驱动 EULA 禁止第三方再分发,且它与内核模块严格 ABI 锁步 | + +⭐⭐ **xim 里已经有这条轴的下半边**:`pkgs/l/libcuda-host-link.lua` 是一个 +**sentinel 包**,只装一个指向 host `libcuda.so.1` 的符号链接,recipe 自己写明: + +> DOES NOT: Redistribute libcuda.so.1. The NVIDIA Driver EULA forbids +> third-party redistribution, and even if it didn't, the userspace lib is in +> strict ABI lockstep with the kernel module — versioning it as an xpkg is +> impossible. + +而且它已经解决了「消费者各自探测 host」的问题:所有 GPU 包读 +`pkginfo.dep_install_dir("libcuda-host-link").."/lib/libcuda.so.1"`, +不各自重实现 ldconfig 探测;宿主库的定位统一走 `libs/hostlib.lua` +(`ldconfig -p` + ELF class 过滤 + first-hit,并有 CI 不变量 +「no recipe enumerates distro library directories」)。 + +#### 由此得到的三条实施要求 + +1. **mcpp 的设备规则包不得从 `PATH` 取 `nvcc`。** 它必须解析到 xim 载荷的路径, + 与现有 `[xlings] deps` 的载荷查找同一条路(`mcpp.build.runner_lookup` 的形状)。 + 本机的 `/usr/bin/nvcc` 只能作为**开发期的回落**,且必须在诊断里说明它来自 host。 +2. **缺的是上半边:xim 里还没有 CUDA 工具链包。** 按 NVIDIA 的 redist 组件划分, + 编译只需要 `cuda_nvcc` + `cuda_cudart`(百 MB 级,不是把整个 toolkit 打进去的 GB 级) + —— 这与 Bazel hermetic CUDA 的按组件取法一致(报告 §2.5)。 +3. ⭐ **这让 §5.2 的宿主编译器配对从「校验」升级为「可满足」**:mcpp 同时提供 + CUDA 版本(xim)与宿主编译器(载荷),所以它可以**选一对兼容的**, + 而不只是在不兼容时报错。本机的现状(载荷 gcc 16.1.0 + host nvcc 12.0, + 上界是 gcc 12)正是「只有校验、无法满足」的样子。 + +### 5.1 本轮只做「可分发」这一档 + +| 载荷 | 用于 | 备注 | +|---|---|---| +| `adaptivecpp` | 形态 B | ⭐ **构建期不需要厂商工具链**(SSCP flow 嵌后端无关 LLVM IR,运行期 JIT),体积最小,是形态 B 成本最低的入口 | +| `oneapi-dpcpp` | 形态 B | 体积大;NVIDIA/AMD 后端还要 Codeplay 插件 | +| `cuda-toolkit` | 形态 A | ⚠️ 数 GB,见 §9 RK-3 | +| `rocm` | 形态 A | ⚠️ 更大 | + +其余(需注册的、编模型的)见 §1.2,不做。 + +### 5.2 ⭐⭐ 宿主编译器配对:一处没人占的位置 + +nvcc 要求宿主编译器版本落在支持表内。报告 §2.2:**CMake 明确不校验** +(#23322 / #24267 / #22202),直接转发 `-ccbin` 让 nvcc 自己报 +`Unsupported gnu version`;xmake 同样不校验。 + +mcpp **自己供给宿主编译器载荷**,所以可以:(1) `-ccbin` 钉到自管载荷; +(2) **在调用 nvcc 之前**校验版本落在支持表内,报错时两边版本都印出来; +(3) 由此让「岛与宿主共享同一份标准库」**由构造成立**。 + +⭐ 第 3 条同时消解了报告 §6 那条结构性障碍(跨边界传 `std::string` 因 +clang/libc++ 撞 nvcc/libstdc++ 而挂)——**它只在两侧标准库不同时成立**。 + +### 5.3 offload 载荷是载荷侧的活 + +带 offload 的 GCC 是一个**独立的载荷变体**(`gcc@16.1.0+offload`), +成本落在 xlings 侧,与引擎并行推进。 + +--- + +## 6. 图与链接 + +### 6.1 岛的产物怎么汇入 + +`.cu` 编出的是**宿主目标文件内嵌设备镜像**(nvcc 的 fatbinary;clang 是 +`.llvm.offloading` 节带 `SHF_EXCLUDE`),就是个普通 `.o`,直接进链接。 +**[已有]** *"Object — its outputs ARE the link edge's inputs"*。 +Metal/Vulkan 的产物走 `Role::Artifact`。 + +### 6.2 device link 只在跨 TU 时需要,且必须声明 + +⚠️ 推断「需不需要」要解析设备代码,mcpp 不做也不该做: + +```toml +[targets.kernels] +device-link = false # 默认。true 时图上多一个 Role::Object 的 action +``` + +⚠️ 它**改变图的形状** ⇒ 必须进指纹,否则改了它 fast path 会 replay 旧图。 + +> 命名用 `device-link` 而非 `rdc`:`rdc` 是 NVIDIA 的词,这件事在 AMD 叫别的、 +> 在 Metal 根本不存在。**引擎的词表不该带厂商口音。** + +### 6.4 不进链接的产物走已有的 `deploy_files` 通道(决定 14) + +Metal 的 `.metallib`、Vulkan 的 `.spv` 是**运行期资源**,不进链接。 +问题是它们怎么随产物走。 + +**[已有]** mcpp 已经有这个通道:`deploy_files = ["bin/widget.dll"]` +(`docs/05-mcpp-toml.md`),经 `runtimeConfig.linkIntent.deployFiles` → +`plan.runtimeDeployFiles` → `ninja_backend`,是当初为 **Windows 无 RPATH 只能把 DLL +拷到 exe 旁**做的。 + +⇒ **规则包把 `Role::Artifact` 的产物登记进 `deploy_files`,不新增产物类别。** +`mcpp pack` 侧也不需要新分类 —— 它已经在处理 deploy 文件。 + +⭐ 这条符合「尽量不加专用东西」:一个 `.metallib` 与一个 `.dll` 对构建系统是同一件事 +——**必须与可执行文件一起送到用户手里的、不参与链接的文件。** + +### 6.3 第一版不产含设备代码的静态库(决定 6) + +报告 §2.2 列了 CMake 的三个未决 issue:#19238、#17586、#21967。 +⇒ **主动避开一个九年没解决的形状。** 要发布走 R3 形态(岛的目标码 + 接缝的模块接口)。 + +--- + +## 7. 身份与分发 + +### 7.1 `accel` 是并列结构化字段(决定 1) + +``` +artifact: + abi : "x86_64-linux-gnu-gcc16-libstdcxx16-c++23" ← 四维字符串,不变 + accel : [ { backend:"cuda", version:"12.8", + archs:["sm_80","sm_90f"], ptx_floor:"90" }, + { backend:"rocm", version:"6.4", archs:["gfx942"] } ] +``` + +**一个比较器,两个存储位置。** `tag_check` 一并比较,不新增第二个比较器 —— +这是本仓库反复踩过的那类缺陷。理由:集合含 `,`/`;`,而现有 tag 是 `-` 拼接、 +从末尾解析三段。**文件名不该承载集合。** + +### 7.2 比较语义:包含 ∪ PTX 下界 + +被接受 ⟺ backend 在集合里,**且** `arch ∈ archs` **或** +(`ptx_floor` 存在且 `arch ≥ ptx_floor`)。 + +| 备选 | 判断 | +|---|---| +| 硬拒 | ❌ 拒掉真能跑的产物,逼回「一颗芯片一份二进制」 | +| 警告后放行 | ❌ **最坏**:问题就是它当前静默,降级成警告等于保持静默 | +| **包含 ∪ PTX 下界** | ✅ | + +⭐ **不是新概念:`standard` 已是 tag 里的非对称下界维**,`tag_check` 已有 +`standard_level()` 的 floor 分支。⭐ AMD 无 PTX ⇒ `ptx_floor` 空 ⇒ 退化成纯包含, +家族覆盖由 ROCm generic target 从集合侧提供。**同一条规则,两家各按自己的机制满足。** + +⚠️ clang 编 CUDA 默认只产 SASS,PTX 要 `--cuda-include-ptx` ⇒ `ptx_floor` 由**实际产物**决定。 + +### 7.3 岛的 C++ 运行时耦合必须核验(决定 7) + +**「接口是 C 形状」不等于「不依赖 C++ 运行时」。** CUDA C++ 就是 C++:岛内部用了 +`std::`/异常/虚函数,目标码就会引用 libstdc++/libc++ 符号。链进用另一套标准库的程序, +就是 memory 里「链接期用 A 运行期加载 B」。 + +⇒ 判据是**目标码里有没有未定义的 C++ 运行时符号**,不是头文件长什么样。 +**[已有]** `mcpp.runtime.elf` 与 `symbol_provision` 已在读符号表,复用。 + +### 7.5 形态 B 的 `accel` 由产物核验得出,节名表由后端提供(决定 13) + +§10.1 指出:形态 B 里 mcpp 无从知道一个目标里到底有没有 kernel, +只知道用户请求了 `-fsycl` ⇒ 光凭声明发 `accel` 会**让产物身份说谎**。 + +⇒ 判据落在产物上:**镜像里有没有该后端的设备段。** + +| 后端 | 节名(示例) | +|---|---| +| CUDA(nvcc) | `.nv_fatbin` | +| clang offload / OpenMP / SYCL | `.llvm.offloading` | +| HIP | `.hip_fatbin` | + +⚠️ 节名逐后端不同,而**引擎不该持有厂商知识**(与 §3.5 规则包持有 flag 拼法同一条原则)。 + +⇒ **引擎读节名,节名表由后端描述符提供。** 引擎侧只需要 +「读 ELF/Mach-O 节名表,判断给定名字是否存在」—— +**[已有]** `mcpp.runtime.elf` 与 `symbol_provision` 已经在做这件事。 + +⭐ 同一个核验顺带答了 §7.3 的问题(岛有没有未定义的 C++ 运行时符号): +**一次读符号/节表,回答两个问题。** + +### 7.4 先修:capability 同名符号边界(决定 4) + +`docs/05-mcpp-toml.md` §2.8.1 实测:绑定**只做选择,不裁剪 link line**; +对「多个 provider 定义同名符号」的能力,两个同时在图里是**缺陷而非可 pin 的歧义**。 + +⚠️⚠️ cuBLAS / rocBLAS / muBLAS 恰好是这一类(白皮书说 muBLAS 刻意与 cuBLAS 同名)。 +⇒ 需在**绑定**这一级拒绝:同一 capability 的多 provider 必须走 +`optional = true` + 互斥 feature。**R6 与 R7 都压在这条上。** + +#### ⚠️⚠️ 这条按原样无法实现,2026-09-05 实施时发现 + +决定 4 是「先修」,而写实现时它立不住: + +- **不能对「一个 capability 有多个 provider 在图里」一律报错。** `docs/05-mcpp-toml.md` + §2.8.1 把 OpenBLAS / MKL 这类**导出不同符号集、按链接选择**的库明确写成**正常用法**, + 一律报错会把它们一起拒掉,而那是一个已发布的、有文档的行为。 +- **也不能只对「符号重叠的」报错**,因为绑定发生在**任何目标文件存在之前**, + 引擎那时读不到符号。`symbol_provision` 能读,但它在链接之后 —— 这正是本条要提前的那一刻。 + +⇒ 缺的是**一个声明**:某个 capability 的 provider 是否互斥,只有能力的定义者知道。 +形如 `provides = ["gpu-blas"]` 之外再声明 `exclusive = true`,或让 capability 本身 +带一个「同名符号」的性质。**这是一个新的 manifest 键与一次设计决定,不是一次实现。** + +因此本轮**没有实现决定 4**,理由不是排期而是它需要先被重新设计。 +⭐ 与之相关的观察:R6 / R7 目前并不真的被它阻塞 —— 互斥 feature 已经能表达这件事 +(`examples` 与索引里的 `compat.*` 都这么用),缺的只是**在用户写错时报错**, +而不是**能不能写对**。 + +--- + +## 8. 生态闭环(决定 10) + +⭐⭐ **引擎认识一个概念 ≠ 有包声明它。** 没有索引侧的包,这套设计对真实开发者 +等于不存在。本节定义「闭环」的具体内容与顺序。 + +### 8.1 四层,每层验证一件事 + +**L0 工具链载荷(xim-pkgindex)** + +| 包 | 验证 | 顺序理由 | +|---|---|---| +| `adaptivecpp` | 形态 B 全链路 | ⭐ **构建期不需要厂商工具链**,体积最小,**能在没有 CUDA 载荷的情况下先把形态 B 跑通** | +| `cuda-toolkit` | 形态 A + `-ccbin` 配对 | 体积风险见 RK-3 | +| `rocm` | 第二家,证明不是 CUDA 特例 | ⭐ **两家跑通才排除「为 CUDA 定制」** | + +**L1 规则包(mcpp-index)** —— `rules-cuda`、`rules-rocm`,各 +`provides = ["mcpp:device-rule="]`。验证 P5 与 `backend` 的 capability 解析。 + +**L2 能力提供者(mcpp-index)** —— `compat.cublas` / `compat.rocblas`, +各 `provides = ["gpu-blas"]`。⚠️ 这两个同时在图里就是 §7.4 的形状, +**所以它们本身就是决定 4 的验证夹具**。 + +**L3 真实的库(mcpp-index)** —— 这一层才是证据。 + +| 包 | 一个包验证几类 | 为什么选它 | +|---|---|---| +| ⭐ **ggml** | R1 + R3 + R5 | **纯 C++、无 Python 层、多后端(CUDA/HIP/Metal/Vulkan/SYCL)、已按后端出预建产物。** 一个包覆盖三类使用者,是单位投入产出最高的 | +| **CUTLASS** | R4 | header-only + 架构门控(`CUTLASS_NVCC_ARCHS` 是按 CUDA 版本门控的白名单),正好压测逐 glob 收窄 | +| **oneMath / oneDNN** | R6 + R8 | UXL 治理的中立库;验证 capability 与形态 B | +| **onnxruntime** | R1 | 已按 EP 分包名,是现成的多变体消费对象 | + +⇒ **首选 ggml。** 理由不是它最流行,是它**一个包同时压到 R1/R3/R5 三类**, +而且它没有 Python 打包层,不会把 wheel 的问题混进来。 + +### 8.1b 要补进生态的包(2026-09-05 追加) + +引擎能力落地之后,生态里必须有**真实流行的**运行时 / 库 / 框架 / SDK, +否则这套能力对开发者等于不存在。下表按依赖顺序排,每一行注明它验证什么。 + +#### xim 侧(工具链与 SDK 载荷) + +| 包 | 内容 | 验证什么 | 依赖 | +|---|---|---|---| +| ⭐ `adaptivecpp` | SYCL,SSCP + **omp 后端** | **形态 B 可在无卡 CI 端到端跑,含 kernel 真执行** | LLVM 载荷 | +| `cuda-nvcc` | `nvcc`/`ptxas`/`nvlink`/`fatbinary`/`cicc` + CRT 头 | 设备工具链不再依赖 host(§5.0b) | —— | +| `cuda-cudart` | CUDA 运行时库与头 | 程序能链、能跑 | `libcuda-host-link`(**已有**) | +| `dpcpp` | intel/llvm 开源版,`configure.py --cuda/--hip` | 形态 B 的 NVIDIA/AMD 后端(无需 Codeplay) | —— | +| `cudnn` | 深度学习原语 | 真实推理框架可用 | `cuda-cudart` | +| `nccl`(后续) | 多卡通信 | 多卡场景 | `cuda-cudart` | + +⚠️ 体积按组件取,不打整个 toolkit(§9 RK-3)。`cuda-nvcc` + `cuda-cudart` 是百 MB 级。 + +#### mcpp-index 侧(规则包、能力包、真实库) + +| 包 | 提供 | 验证什么 | +|---|---|---| +| `rules-cuda` | `mcpp:device-rule=cuda` | P5 规则包;`backend` 经 capability 解析 | +| `compat.cudart` | `gpu-runtime` | capability 绑定 | +| `compat.cublas` | `gpu-blas` | ⚠️ 与 `compat.rocblas` **同名符号** ⇒ 它们本身就是决定 4 的验证夹具 | +| `compat.cccl` | header-only(Thrust/CUB/libcu++) | 最小可用的 GPU 算法库;无二进制分发问题 | +| `compat.cutlass` | header-only,按 CUDA 版本门控架构白名单 | **逐 glob 架构收窄(R4)的压测对象** | +| ⭐⭐ `llama.cpp-m` 加 CUDA 后端 | —— | **R1 + R3 + R5 三类;而且它已经在 `mcpplibs/` 里** | +| `compat.onnxruntime`(后续) | 按 EP 分变体 | R1 的多变体选择,真实体量 | +| `compat.opencv` + CUDA(后续) | —— | 真实的混合工程 | + +⭐⭐ **`llama.cpp-m` 是首选验证对象,理由不是它流行,是它一个包同时压到三类使用者**, +而且它没有 Python 打包层,不会把 wheel 的问题混进来。它已经存在于生态中, +所以这一步是「给它加一个后端」,不是「从零收录一个大工程」。 + +⭐ `compat.cccl` 与 `compat.cutlass` 都是 header-only,**没有二进制分发与体积问题**, +所以它们是 §11.1 说的「源码分发这条最短路径」上最先能落地的两个真实库。 + +### 8.2 ⚠️⚠️ 发布顺序是一个环,必须拆开 + +memory 记过两条硬的:「消费者先发布,索引 `latest` 才能动」、 +「⚠️⚠️ PR 绿 ≠ 合入后 main 绿,判据是 `origin/main` HEAD SHA 上的 run」。 + +这里的环是: + +``` +引擎支持 accel → 索引包才能发 accel → 才能做真实验证 → 才知道 accel 形状对不对 + ↑ │ + └────────────────────────────────────────────────────────────────────┘ +``` + +拆法:**用 path 依赖与本地包在索引之外先闭环一次**,再发布。 +即阶段顺序必须是「本地闭环 → 引擎发布 → 索引发布 → 沙箱验证」,四步不能并。 + +⚠️ 而 §9 RK-2 说明:**`accel` 字段进已发布描述符之前,必须先确认旧客户端读它不会失败。** + +### 8.3 ⭐ 真实验证:分清哪些判据不需要 GPU + +这是应对 RK-1 的关键,也是本节最有价值的一条。 + +**不需要 GPU 就能验的(占绝大多数)**: + +- `tag_check` 的拒绝与诊断文本 +- 架构收窄成空集/非子集的报错 +- 宿主编译器配对不合时**在 nvcc 之前**报错 +- `.cu` 不出现在 P1689 结果里;改 `.cuh` 触发重编 +- `cfg(any(accelerator=…))` 的合并与未知键诊断 +- 同名符号的 provider 在**绑定期**被拒 +- **编译与链接本身** —— 在无卡机器上编 `sm_90` 的代码完全正常 + +**必须有 GPU 才能验的(少数,要隔离)**: + +- kernel 真的跑起来 +- `--accel native` 的探测 +- 运行期多后端派发(R5) + +⚠️⚠️ memory:「`# requires: llvm` 的 e2e 从未在 CI 跑过 —— 两 shard 都无 llvm 且 +skip 退 0;⭐**守卫住在 job 里**」。 +⇒ **需要 GPU 的用例不得与其余 e2e 混在同一套件里用 `# requires:` 跳过**, +否则它们会永远绿而从未运行。它们必须住在**一个独立的、可见其存在与否的 job** 里; +该 job 不存在时,状态是「未覆盖」而不是「通过」。 + +### 8.5 ⭐⭐ 没有 GPU,能验到哪一步 + +「有没有模拟 GPU」这个问题的答案分三档,而它**改变了排期的理由**。 + +**第一档 —— 编译期判据,完全不需要 GPU(占绝大多数)** + +`tag_check` 的拒绝与诊断、架构收窄成空集/非子集的报错、宿主编译器配对提前报错、 +`.cu` 不进 P1689、改 `.cuh` 触发重编、`cfg` 合并与未知键诊断、绑定期同名符号拒绝、 +以及**编译与链接本身** —— 在无卡机器上编 `sm_90` 的代码完全正常。 + +**第二档 —— kernel 可以在 CPU 上真跑(这一档我原来漏了)** + +| 后端 | CPU 上跑 kernel 的路径 | 证据强度 | +|---|---|---| +| **SYCL / AdaptiveCpp** | `--acpp-targets=omp`(`omp.library-only` / `omp.accelerated`);其后端列表含 **host CPU (LLVM)** | ⭐ 已从 AdaptiveCpp 官方 `doc/compilation.md` 核实 | +| SYCL / Intel DPC++ | OpenCL CPU runtime 提供一个 CPU device | 未在本轮核实 | +| Vulkan compute | Mesa **lavapipe** / Google **SwiftShader**(纯软件 Vulkan 实现) | 未在本轮核实 | +| OpenCL | **PoCL** | 未在本轮核实 | + +**第三档 —— 没有生产可用的模拟路径** + +- **CUDA**:`-deviceemu` 早已从工具链移除;GPGPU-Sim / Accel-Sim 是学术模拟器, + 速度与版本跟进都不适合 CI。⇒ **CUDA kernel 的执行只能上真卡。** +- **HIP**:有 HIP-CPU(header-only),覆盖面受限。 +- **Metal**:没有 CPU 实现。 + +#### ⭐⭐⭐ 由此改写 AdaptiveCpp 优先的理由 + +v4 说先做 AdaptiveCpp 是因为「载荷体积最小、构建期不需要厂商工具链」。 +这是对的但不是最重要的。真正的理由是: + +> **形态 B 经 AdaptiveCpp 的 omp 后端,可以在没有任何 GPU 的 CI 里 +> 端到端跑完 —— 包括 kernel 真的执行。** + +也就是说,**RK-1(CI 无 GPU)对形态 B 几乎不成立**;它只对形态 A 的 +「kernel 真的跑起来」这一条成立。⇒ 排期上先形态 B,不只是省钱,是**先拿到一条 +可以完整自证的链路**,再去做只能部分自证的形态 A。 + +### 8.4 生态 CI 从 PR 分支现场构建 + +memory:「生态 CI 从 mcpp 的 PR 分支现场构建(`MCPP_SOURCE_REF`);⭐⭐按 step 穷举 +不是按仓库;⚠️没有任何 CI 解析已发布的包 ⇒ 沙箱那步唯一」。 + +⇒ 索引侧的加速器包必须在 **mcpp PR 分支的引擎**上验证,否则验的是旧引擎; +而「已发布物真的能用」只有沙箱那一步能答。 + +--- + +## 9. 风险 + +按「做错了谁能自救」排序。 + +### RK-1 ⚠️⚠️ CI 没有 GPU,而这套设计的目的正是防一个运行期错误 + +**这是最大的风险。** 本设计存在的理由是把 +`no kernel image is available for execution on the device` 从运行期提前到构建期, +而项目**没有 GPU 可以证明这个提前真的发生了**。 + +历史先例就在本仓库:`# requires: llvm` 的 e2e 因两个 shard 都无 llvm、skip 退 0, +**从未在 CI 跑过而一直是绿的**。 + +**缓解,分两层**: + +1. §8.3 —— 绝大多数判据不需要 GPU;需要的那少数**隔离到独立 job**, + job 不存在时读数是「未覆盖」而非「通过」。 +2. ⭐⭐ §8.5 —— **形态 B 经 AdaptiveCpp 的 omp 后端可以在 CPU 上真跑 kernel**, + 于是这条风险**对形态 B 几乎不成立**。先做形态 B 就是先拿到一条能完整自证的链路。 + +**残余风险(缩小后)**:仅剩**形态 A(CUDA/HIP)的 kernel 执行**与 +R5 的运行期多后端派发,无生产可用的模拟路径,只能上真卡或标为未覆盖。 + +### RK-2 ⚠️ 旧客户端会把 GPU 产物当成「无约束」而接受(等级已下调,但形状变了) + +v4 把这条列为 ⚠️⚠️「旧 mcpp 可能整份 manifest 加载失败」。**读代码后这条不成立。** + +`modules/manifest/src/xpkg.cppm`(已发布描述符的读取器)注释逐字: + +> the mcpp-segment key vocabulary is a CLOSED whitelist (the parse loop's else-if +> chain). An unrecognised key is collected into `Manifest::xpkgUnknownKeys` and +> **silently skipped** + +并且有 `closest_known_xpkg_key` 把未知键映射回最可能的拼写做提示。 +⇒ **旧客户端读到 `accel` 会跳过它,不会失败。**「新功能要新 mcpp」这个前提是成立的。 + +⚠️ **但残余风险换了个形状,而且正是我们要防的那个:** +旧客户端跳过 `accel` 之后,会把一个 `sm_90` 的产物当作**没有加速器约束**而接受, +在 `sm_86` 的机器上装上去 —— **`no kernel image` 又回来了,只是发生在旧客户端上。** + +⭐ **缓解不需要任何新机制,只需要产物顺序。** 消费循环是 first-match-wins: + +```cpp +auto bad = tag_check(*published, in.current); +if (bad.empty()) { accepted = true; break; } +``` + +⇒ **把 CPU-only 变体放在 `artifacts` 列表的第一个。** +旧客户端(不看 `accel`)拿到它 —— 安全、能跑、只是慢; +新客户端逐个评估 `accel`,跳过不匹配的,挑到真正合适的那个。 + +⇒ 这条降级为「发布约定」,写进 §11.2,并由 `mcpp pack` 保证顺序,不靠人记。 + +### RK-3 ⚠️ 载荷体积 + +ConanCenter 当年的第二条理由逐字是 *"Binaries for CUDA are huge (several Gb)"*, +这条到今天仍然成立。memory 另记过「GitCode 镜像上传限速 + 载荷瘦身 —— **主因其实是没 +strip**」。 + +⚠️⚠️ 另有一条运维上的硬约束(memory):**GitCode 资产不可替换、不可删除 ⇒ +同号不同字节做不到。** 一个发错的多 GB 加速器产物**永远改不回来**,只能 bump 版本。 +产物越大、变体越多,这条的代价越高。 + +**缓解**:(1) 只打可再分发子集,不打整个 toolkit;(2) strip; +(3) ⭐ **先做 AdaptiveCpp**(§8.1)—— 它构建期不需要厂商工具链,可以在不解决体积问题 +的前提下先把形态 B 的全链路跑通;(4) 大变体发布前先在沙箱验一遍(§8.4), +因为发出去就改不回来了。 + +### RK-4 ⚠️ 变体矩阵由谁构建 —— mcpp-index 的 CI 也没有 GPU + +一个包 × N 后端 × M 架构家族的产物要有人编出来。**但编译不需要 GPU**(RK-1 的划分), +所以这条比看起来轻:CI 能编,只是不能跑。 +**残余风险**:编出来的产物没有被运行验证过就进了索引。 + +⭐ **但不加标注(决定 16)。** 理由:**mcpp 今天已经在发布交叉编译的产物 —— +一个在 x86 上编出来的 aarch64 产物,同样没有在目标机器上跑过。** +GPU 产物是同一形状,不是新问题。为它单独加一个字段,等于承认交叉编译产物有两个等级, +而这个区分从来没有被建立过,也没有读者。 +⇒ 与其加字段,不如把「哪些后端能在 CI 里真跑」这件事(§8.5)做成排期依据。 + +### RK-5 法律与许可 + +CUDA EULA 的 Attachment A 允许随应用再分发指定组件,NVIDIA 也发布 +`redistrib_.json` 机器可读清单 —— 但清单里只写 `"license": "CUDA Toolkit"`, +**如何解析成实际许可文本仍是开放问题**(NVIDIA/build-system-archive-import-examples#3)。 + +**缓解**:载荷描述符逐条记录来源 URL 与 sha256(与 Bazel hermetic CUDA 同法), +许可字段照抄清单原文不做解释。 + +### RK-6 ⚠️ 上游包的身份规则 + +memory:「上游包的 namespace 与版本都用上游的;**只有内容自己写的才用 mcpplibs**; +⭐『对齐上游』要逐字节 cmp;⚠️ licence 是**集合**,写错一个比不写更坏」。 + +⇒ `cuda-toolkit`、`rocm`、`adaptivecpp` 用上游身份; +`rules-cuda`、`compat.cublas` 是我们写的内容,用 mcpplibs。 + +### RK-7 形态 B × modules 未验证(§2.3) + +若实测 icpx/AdaptiveCpp 不支持含 SYCL kernel 的 module interface, +形态 B 的目标必须整体退回 header 模式 —— 这会削弱「mcpp 是 modules-first」的卖点。 +**缓解**:阶段 0b 先测,再决定形态 B 的排期。 + +### RK-8 ⚠️ 与 CMake 的易用性差距是设计取舍,不会消失 + +「我有一个 `.cu`」在 CMake 是 `project(x LANGUAGES CXX CUDA)` 加一行架构; +在 mcpp 要理解 device target、规则包、接缝三个概念。 +**根因是 mcpp 用接缝换来了后端可替换(R5/R7),CMake 没有这个能力所以可以更简单。** +**缓解**:三行最小形态(§10.4)、`mcpp new --template cuda`、文档第一例必须是三行那个。 + +--- + +## 10. 自我 review + +按四个维度,**只写查出来的问题**。 + +### 10.1 架构设计 + +**✅ v4 的两处最弱本轮都收口了**:形态 B 的身份改为**产物核验**(决定 13,§7.5), +`backend` 的两趟归属写死为「请求在第一趟、绑定在第二趟、缝上有 satisfied 记录」 +(决定 12,§3.5),后者直接套用了 mcpp 已有的「请求侧与答案侧各留拼写」经验。 + +**⚠️ 现在最弱的一处:两种形态是两套心智模型,而这是真实的,不是包装问题。** +形态 A 加一个 target,形态 B 改 `[toolchain]`。字段名统一了(决定 15)、身份层统一了 +(P3/P4 两边都用),但**「我要用 GPU」这句话在两种形态下要做的事仍然不一样**。 +我找不到一个不撒谎的统一写法 —— 底层模型确实不同,**强行统一会是假的统一**。 +⇒ 处理办法只能是文档把两条路分开讲清楚,而不是假装它们是一条。 + +### 10.2 简洁优雅 + +**✅** 相对 v2 少了两个键(`archs-intersect` 合并、`accelerator-any` 去掉); +`backend` 走 capability;工具链沿用既有 `family@version` 而不是新发明一套。 + +**⚠️** 形态 A 与 B 的 manifest 写法完全不同(一个加 target,一个改 `[toolchain]`)。 +这反映底层模型真的不同,**强行统一会是假的统一** —— 但用户要学两套。 + +**✅ 决定 15 已把字段名统一成 `archs`**(v4 自查发现的不一致)。 + +**⭐ 本轮新增的一项自查:有没有「乱加专用的东西」。** +把全部新增的用户可见面逐条列出来数一遍: + +| 类别 | 新增 | 数量 | +|---|---|---| +| manifest 键 | `kind="device"`、`backend`、`archs`、`device-link`、`[package] accelerators`、`[build] accel` | 6 | +| cfg | `accelerator` 一个 layer key(**复用已有的 `any`/`all`,不加谓词**) | 1 | +| CLI | `--accel` / `--no-accel` | 2 | +| 描述符 | `accel` 字段 | 1 | +| **新命令(`mcpp `)** | —— | **0** | +| **新机制** | —— | **0** | + +**原样复用、一处未改的**:`BuildAction::Role`、`tag_check` + 消费循环、capability 绑定、 +`deploy_files`(决定 14)、`mcpp.runtime.elf`(决定 13)、`CommandDialect`、 +`[package] platforms` 的形状(§11.1)、`[toolchain] family@version`(决定 8)、 +规则包、11 字段指纹。 + +⭐ `--accel` / `--no-accel` 不是新花样:它与 `[build] accel` 的关系, +和 `--target` ↔ `[toolchain]`、`--features` ↔ `[features]` 完全同形 +—— **manifest 声明 + 命令行覆盖**,mcpp 已有的模式。 + +### 10.3 兼容性 + +**✅** `.cu` 进分类表但**不进默认 glob**,`builtin_extension_table()` 的 +`.ixx` 判例逐字适用:*"a break its author cannot fix, because the tarball for that +version has already shipped."* + +**✅ RK-2 由代码给出了答案**:已发布描述符的读取器对未知键是 +*"collected into `Manifest::xpkgUnknownKeys` and **silently skipped**"* ⇒ +旧客户端不会加载失败。「新功能要新 mcpp」这个前提成立。 + +**⚠️ 但残余风险换了形状**:旧客户端跳过 `accel` 后会把 `sm_90` 产物当成无约束而接受。 +缓解是**产物顺序(CPU-only 排第一)**,零新机制 —— 但它必须由 `mcpp pack` 保证, +**不能靠发布者记住**,否则就是一条没有执行者的约定。 + +### 10.4 易用性 + +**✅** R1(数量最大)零改动;错误信息本身就是产品。最小形态可以是三行: + +```toml +[targets.kernels] +kind = "device" +sources = ["src/*.cu"] +``` + +(`backend` 在只有一个 device 规则包时自动绑;`archs` 缺省 ⇒ **报错并列出该 backend +的常用取值**,不猜 —— 理由是报告 §2.2 的 CMake `native` 陷阱,根因就是把 +「本机有什么」当成「要编成什么」。) + +**⚠️** RK-8:与 CMake 的差距是取舍不是缺陷,但要说出来而不是粉饰。 + +### 10.5 ⭐ 生态闭环(本轮新增的维度) + +**⚠️ 最弱的一处:L3 的「真实可用」目前只有一个候选扛得住。** +ggml 之外,CUTLASS 是 header-only(不验证分发)、oneMath 依赖形态 B(依赖 RK-7)、 +onnxruntime 体量太大不适合首发。 +⇒ **整个生态验证在第一阶段实际压在 ggml 一个包上**,它若卡住(例如它的 +多后端构建对 mcpp 的图有超出预期的要求),闭环就没有第二条腿。 +**缓解**:L3 之前先用一个**自建的最小多后端库**做夹具,把风险与 ggml 解耦。 + +**⭐ 本轮新增:§8.5 让「真实可用」的第一步比预想的近。** +形态 B 经 AdaptiveCpp 的 omp 后端可以在无卡 CI 里**端到端跑完并真的执行 kernel**, +而 §11.1 说明**源码分发不受体积与「谁来编」两条风险影响**。 +⇒ **「AdaptiveCpp + 源码分发」是一条能完整自证、且没有运维风险的最短路径**, +应当作为第一条打通的链路;CUDA + 二进制分发排在其后。 + +--- + +## 11. 源码分发与二进制分发 + +加速器把这两条路的差异放大了,所以要分开说清楚。 + +### 11.1 源码分发 —— 消费者的 `archs` 说了算 + +源码包被消费时,**是消费者在编**,所以 `archs` / `backend` 由**消费者**决定, +包本身不含任何架构。那么包该声明什么? + +⭐ **复用 `[package] platforms` 的形状**:mcpp 已有 +`platforms = ["linux","macos","windows"]`,文档说它是 +*"the platforms the package supports (a CI matrix hint, shown via `mcpp why`)"* +—— **一个声明,不是强制。** + +```toml +[package] +platforms = ["linux", "windows"] +accelerators = ["cuda", "rocm"] # [提议] 同一形状:支持面声明,非强制 +``` + +⇒ **两个字段语义严格对称**:`platforms` / `accelerators` 说「我支持什么」(声明), +`accel` 说「这个二进制里有什么」(实测)。**声明面与实测面各有各的字段,不混。** + +三条后果: + +1. 源码分发时消费者**必须自己有工具链** ⇒ 走 xlings 载荷自动安装,与 gcc/llvm 同一条路。 +2. ⚠️ **编译成本落到消费者头上。** CUTLASS 规模的模板实例化,一个错的 `archs` + 可能是一小时。这是决定 11(缺省报错不猜)的第三条理由。 +3. ⭐ 源码分发**不受 RK-3(体积)与 RK-4(谁来编)影响** —— + 它是加速器支持成本最低的一条路,也应该是**第一条打通的路**。 + +### 11.2 二进制分发 —— 粒度、顺序、以及不可撤回 + +**粒度:一个产物一个家族目标。** +不是每颗芯片一份(组合爆炸,报告 §3.5:Spack 旗舰 buildcache 里 GPU 二进制为零), +也不是全塞一份(体积)。家族目标(`sm_90f` / `gfx10-3-generic`)是 2025 年才有的 +中间选项,正好是这个粒度(报告 §1.3)。 + +**顺序:CPU-only 变体必须排第一。** 见 RK-2 —— +消费循环是 first-match-wins,旧客户端不看 `accel`,谁在前面就拿谁。 +⇒ **`mcpp pack` 保证这个顺序,不靠发布者记住。** + +**不可撤回:** ⚠️⚠️ GitCode 资产不可替换不可删除 ⇒ 多 GB 的产物发错只能 bump 版本。 +⇒ 大变体发布前必须过沙箱(§8.4)。 + +### 11.3 ⭐ 实际形态是混合的,设计不该假设二选一 + +生态里真实发生的事: + +- **kernel 走源码**(消费者按自己的卡编), +- **厂商库走二进制**(cuBLAS/cuDNN 本来就只有二进制), +- **框架走二进制**(ggml/onnxruntime 按后端出预建产物)。 + +⇒ 三者在同一个依赖图里共存。这对设计的要求是: +**`accel` 的「缺席=不受约束」语义必须在两条路上都成立** —— +源码包没有 `accel`(它还没被编出来),CPU-only 二进制也没有 `accel`(它不受约束), +两者对 `tag_check` 是同一个读数,而这**正好是对的**:两者确实都不施加加速器约束。 + +⭐ 这条是「空维=不受约束」这个已有语义的一次意外红利 —— +它让源码包与 CPU 包在身份系统里天然同形,不需要第三种状态。 + +--- + +## 12. 分阶段与判据 + +判据取**只有做对了才会出现的读数**,不取「构建通过」。标注 `[需GPU]` 的按 §8.3 隔离。 + +| 阶段 | 内容 | 判据 | +|---|---|---| +> ⭐ **排期主线(§8.5 + §11.1 的结论):先「AdaptiveCpp + 源码分发」,再「CUDA + 二进制分发」。** +> 前者能在无卡 CI 里端到端自证(kernel 真跑),且不受体积与「谁来编」两条运维风险影响; +> 后者两样都占。下表的 L0/8 之所以排在 CUDA 相关项之前,就是这个理由。 + +| **0a** | 兼容性探针(RK-2) | 旧版 mcpp 读手工加了 `accel` 的描述符:**忽略而非失败**(代码已表明是 skip,此处是复核) | +| **0b** | 形态 B 探针(RK-7) | 实测 icpx / AdaptiveCpp 能否编含 SYCL kernel 的 module interface;以及是否需要新 `CompilerId` | +| **0c** | 规则包原型(零引擎改动) | 一个带 kernel 的工程能 `mcpp build` 出可运行二进制 `[需GPU]` | +| **1** | capability 同名符号边界(决定 4) | 两个 `provides=["gpu-blas"]` 同时在图里 ⇒ **绑定期报错**,非链接期重复符号 | +| **2** | `accel` 字段 + `tag_check` | 同一包发两变体,消费者按声明架构选中不同的那个;⚠️ 两变体须在其余各维完全相同 | +| **3** | `SourceKind::Device` + `.cuh` 进 header 轴 | `.cu` 不在 P1689 结果里;**只改一个 `.cuh` 后重建,kernel 对象真被重编**(判据落产物内容) | +| **4** | `accelerator` 多值 layer + 组合子语义 | `cfg(any(accelerator="rocm"))` 在 cuda-only 构建下不合并;拼写错被 `unknown_tokens()` 报出 | +| **5** | 逐 glob 收窄(决定 3) | 空集 ⇒ 报错;非子集 ⇒ 报错 | +| **6** | `-ccbin` 钉定 + 配对校验 | 配一个超出支持表的宿主编译器,**mcpp 在 nvcc 之前报错**,消息含两边版本号 | +| **7** | device link | 跨 TU `__device__` 调用能链上;改 `device-link` 触发重新 prepare | +| **L0** | `adaptivecpp` 载荷 | ⭐ 无 CUDA 载荷、**无 GPU** 的机器上,形态 B 全链路跑通,**且 kernel 真的执行**(`--acpp-targets=omp`) | +| **L1/L2** | `rules-cuda` + `compat.cublas`/`compat.rocblas` | 两个 provider 同时在图 ⇒ 绑定期报错(与阶段 1 同一判据,但这次是**真实的包**) | +| **L3a** | 自建最小多后端夹具(**源码分发**) | 同一份源码,消费者按自己的 `archs` 编出可用产物;不涉及二进制发布 | +| **L3b** | ggml(**二进制分发**) | 同一份源码,`--features cuda` 与 `--features rocm` 产出 `accel` 不同的两个产物;CPU-only 变体排第一(§11.2) | +| **8** | 形态 B(SYCL/OpenMP) | 换 `[toolchain]` 后产物里出现设备段;offload 载荷有 `nvptx-none` 的 `libgomp` 插件 | + +#### 2026.9.5.1 实际落地了哪些阶段 + +| 阶段 | 状态 | 证据 | +|---|---|---| +| 0a 兼容性探针 | ✅ | `xpkg.cppm` 读取器对未知键 skip;`mcpp.toml` 的 `runtime.artifacts` 是封闭白名单,已加 `accel` | +| 0b 形态 B 探针 | ❌ 未做 | 本机无 SYCL 工具链,做了无法真实验证 | +| 0c 规则包原型 | ✅ | `examples/09-cuda-kernel`,RTX 4080 上 `mcpp run` 输出 `12 24 36 48` | +| 1 capability 边界 | ❌ **未做,且需重新设计**(见 §7.4 的补注) | +| 2 `accel` + `tag_check` | ✅ | 单测 12 条 + e2e 600;真实拒绝消息见 §14 | +| 3 `SourceKind::Device` | ✅ | 单测 6 条;`.cuh` 进 header 轴 | +| 4 `accelerator` 多值 layer | ✅ | 单测 7 条;语义按 §3.4 修正后的一行 | +| 5 逐 glob 收窄 | ❌ 未做 | 依赖 device target(阶段外),本轮 device 编译走规则包 | +| 6 `-ccbin` 配对 | ⚠️ 部分 | 引擎侧读上界并在 `mcpp self doctor` 报告;**钉定**发生在规则包里,不在引擎里 | +| 7 device link | ❌ 未做 | 依赖 device target | +| L1/L2 索引包 | ⚠️ 部分 | `compat.cuda-runtime` 已合入并发布并沙箱验证;`rules-cuda` 目前是 `examples/` 里的 path 包,尚未收录进索引 | +| L0 / L3 / 8 | ❌ 未做 | 见上 | + +⚠️ **决定 3(逐 glob 收窄第一版就做)与决定 4(capability 边界先修)都没有落地。** +前者依赖 device target 这一原语,而本轮 device 编译走的是规则包路线(设计 §12 阶段 0c +本来就排在引擎 device target 之前);后者见 §7.4 的补注 —— 它需要先被重新设计。 +两条都不是「忘了」,但都与拍板时的意图有出入,记在这里而不是让阶段表替它们含糊过去。 + +⚠️ **0a 是前置不是并行**:它若答「失败」,阶段 2 的形状要改。 + +--- + +## 13. 仍未决 + +v4 的六条已全部拍板(决定 11–16)。真正剩下的: + +1. **`archs = ["native"]` 是否提供。** 决定 11 拒绝了「缺省时猜」,但没决定 + 「显式写 native 时探测」要不要给。倾向给,且**只在本地,CI 里用它应当告警** —— + 否则就是把 CMake 那个陷阱换个写法搬进来。 +2. **形态 B 的 `archs` 取值空间与形态 A 不同**(`spir64` / `nvptx-none` 是 triple 形状, + `sm_90f` 是 arch 形状)。字段名统一了(决定 15),但**取值该不该也统一**, + 还是由 backend 各自定义值域?倾向后者,但需要写进 §7.1 的 `accel` 结构说明。 +3. **`mcpp pack` 一次能不能产多个 `accel` 变体**,还是每个变体一次调用。 + 影响 RK-4(谁来编)与 §11.2 的顺序保证。 +4. **规则包提供节名表的接口形状**(决定 13)—— 是规则包的模块导出一个函数, + 还是描述符里的一张静态表?后者更适合「引擎不加载规则包也要能核验产物」。 + +--- + +## 14. 实施后的自我 review(2026-09-05,PR #559 / mcpp-index #346) + +本节只记**实施推翻了设计的地方**,以及生态级的观察。通过的不记。 + +### 14.1 设计被实施推翻的五处 + +| # | 设计说 | 实测 | 影响 | +|---|---|---|---| +| 1 | `any(accelerator=…)` 表示成员判定,裸键表示集合相等(§3.4 三行语义) | **错的。** 让组合子改变操作数含义会使 `all(a="cuda", a="rocm")` 变成不可满足而不是「两者都启用」 | 改为**处处成员判定**,组合子保持纯布尔。**用户少学一条规则**,`cfg(accelerator="cuda")` 直接可用 | +| 2 | 「xpkg 已确认 skip 未知键 ⇒ RK-2 等级下调」 | **只对一条路成立。** 已发布描述符的读取器确实 skip;而**打包后的 `mcpp.toml` 的 `runtime.artifacts` 是封闭白名单,未知键直接报错** | RK-2 在 mcpp.toml 这条路上是真的。两个读取器两种策略,设计只查了一个 | +| 3 | 形态 B 需要一个新的工具链家族(§5.0 倾向) | **不需要。** DPC++ 就是 clang:同 driver、同 GNU flag、同 BMI。`Family` 封闭枚举一处都不用动 | 形态 B 的工程量从「加家族」降到「加载荷变体 + 能力位」 | +| 4 | nvc++ 提供 `-stdpar`,看起来属于形态 B | **属于形态 A。** HPC SDK 自述主标准是 C++17,无 modules ⇒ 当不了 `import std` 工程的宿主编译器 | 它只能作设备编译器,与 nvcc 同类 | +| 5 | 宿主编译器上界内置一张按 CUDA 版本的表(§5.2 备选) | **应当读 `crt/host_config.h`。** 表是某一个 release 的拷贝,下一版就静默过时 | 一个 mcpp 从未见过的 toolkit 也能作答;解析不了的头文件不产生断言 | + +⭐ 第 1 条值得单独说:它是**实施让设计变简单**的一次,而不是变复杂。 + +### 14.2 生态强制执行了设计原则,而不是相反 + +三次实测里,**生态自己拦下了我想走的近路**: + +1. **链宿主 `libcudart` 被拒。** mcpp 的运行期闭包校验报 + *"Its PT_INTERP is a private loader, so the host's /usr/lib is NOT consulted"*, + 并指向 `xlings install `。**这不是我设计的,是引擎已有的。** + §5.0b 的「少依赖 host」不是一条建议,是被执行的规则。 +2. **`allow_host_libs = true` 不够。** 闭包校验比它更强。只能改成静态链接可再分发的一半, + 把宿主依赖收敛到 `libcuda.so.1` 这一个不可消除的点。 +3. **`[xlings] deps` 只对 ROOT 工程 materialize。** 索引里 `riscv-virt-rt` 的注释已经写着 + 这条,并给出正解:包级的安装边是 `xpm..deps`。 + +⇒ **生态的既有约束比我的设计更严,而且它们是对的。** + +### 14.3 ⚠️⚠️ 我自己犯的两个错,都是本文档反复警告过的形状 + +**其一,重复了一处被明令禁止重复的探测。** `compat.cuda-runtime` 的第一版自己抄了一份 +宿主 `libcuda` 探测,而 xim 的 sentinel recipe 原文写着: + +> Single source of truth for "where is host libcuda" → all GPU xpkgs read from +> `pkginfo.dep_install_dir(...)` and **don't reimplement ldconfig probing each** + +xim 的 `hostlib.lua` 还记着这条规则的来历:**四处各自探测,三处是错的**, +每一处都是同一个「假设目录布局」的错误。我写了第五处。 +**是 review 抓出来的,不是我自己发现的。** + +**其二,凭印象写下了一条「为什么需要」的理由,而它是错的。** +第一版还抓了 `libnvidia-ptxjitcompiler`,PR 正文写「没有它前向兼容会静默失效」。 +实测:只编 `compute_80`、只让 `libcuda.so.1` 可见,在 sm_89 上 JIT 成功(`jit result=99`)。 +驱动通过自己的路径加载兄弟库。**三个 pattern 全是多余的。** + +⭐ 这正是 memory `reasons-written-from-memory-kill-good-fixes` 的镜像: +凭印象写的理由既能否掉正确修法,也能**留下多余实现**,而**结论会被复查,理由不会**。 + +### 14.4 生态级:一个跨工具的缺陷 + +`xlings install` 会把**由 mcpp 安装的**裸名 shim 从 subos 的 `bin/` 里剪掉, +包括 `mcpp` 自己,即使那次安装是 already-installed 的 no-op。 +已受控复现(13 条 → 6 条,删掉 7 条),已提 openxlings/xlings#582。 + +⭐ 形状是:**两个工具写同一个 shim 目录,其中一个把另一个的条目当作陈旧项。** +这与本文档 §1 的主题同构 —— 一个字段两个写者,而其中一个不知道另一个存在。 + +### 14.5 本轮的覆盖与未覆盖 + +**已落地并真机验证**(RTX 4080 / CUDA 12.0 / 驱动 550.144.03): +`SourceKind::Device`、`accelerator` 多值 layer、`accel` 身份维与 `tag_check`、 +描述符与构建请求两端接线、宿主/设备编译器配对(`mcpp self doctor`)、 +`--accel` / `--no-accel`、`[package] accelerators`、docs 20 章中英双份、e2e 600、 +`examples/09-cuda-kernel`(`mcpp run` 输出 `12 24 36 48`)、 +`compat.cuda-runtime`(farm 单条链向 sentinel)。 + +**未做,且理由是「做了无法真实验证」**:形态 B(本机无 SYCL 工具链)、 +device link / RDC、含设备代码的静态库、Metal(本机无 macOS)、 +xim 的 CUDA 工具链载荷。每一项在 §12 都有独立判据。 + +### 14.6 §8.2 的环被实际走了一遍,并且闭合了 + +设计说跨仓库顺序是硬的,而这一轮把它走完了: + +1. 引擎侧改动进 PR #559(未发布); +2. `compat.cuda-runtime` 只用现有键,因此**不必等引擎发布**,独立进 mcpp-index #346; +3. #346 合入后,**判据是 `Publish Index Artifact` 在合入的那个 commit 上绿** + —— 不是「PR 绿」,也不是「main 有这个文件」。实测该 workflow 在 `a6f625e3` 上 + completed/success 之后,`mcpp index update` 才能解析到它; +4. 于是示例把 `LD_LIBRARY_PATH` 的绕法换成一条普通依赖,`mcpp run` 无环境变量直接输出 + `12 24 36 48`。 + +⭐ 第 3 步是 memory `index-stale-but-marker-fresh` 记的第五层: +**xlings 消费的是 `artifact:`**,所以「合入了」与「能用了」之间还隔着一个 workflow。 +本轮没有踩到它,因为事先按那条记忆等了发布再验证。 + +⚠️ 仍未闭合的是另一件事:**引擎侧的新键要等 #559 发布之后,索引里的包才能使用**。 +`compat.cuda-runtime` 恰好一个新键都没用到,所以这一轮绕开了; +下一个用到 `accel` 字段的索引包不会这么幸运。 diff --git a/.agents/docs/2026-09-05-accelerator-support-implementation-plan.md b/.agents/docs/2026-09-05-accelerator-support-implementation-plan.md new file mode 100644 index 000000000..5c209e399 --- /dev/null +++ b/.agents/docs/2026-09-05-accelerator-support-implementation-plan.md @@ -0,0 +1,303 @@ +# 加速器支持 实施计划 + +> **For agentic workers:** 用 superpowers:subagent-driven-development 或 +> superpowers:executing-plans 逐任务实施。步骤用 `- [ ]` 跟踪。 + +**Goal:** 让 mcpp 能声明、构建、发布并正确选择加速器产物,并在 mcpp-index / xlings +上打通一条真实可用的生态链路。 + +**Architecture:** 五个原语(device target / 接缝模块 / `accel` 身份字段 / +`accelerator` cfg layer key / 规则包),其中三个复用已有机制。引擎持有身份、指纹、 +架构集合与工具链配对;厂商 flag 拼法由规则包持有。 + +**Spec:** [`2026-09-05-accelerator-support-design.md`](2026-09-05-accelerator-support-design.md) + +**Tech Stack:** C++23 modules,mcpp 自举构建,ninja 后端,xmake 无关。 + +## Global Constraints + +- 基线一律 `origin/main`(d18cd8b,2026-09-04)。worktree: + `/home/speak/workspace/github/mcpp-community/mcpp-accel`,分支 `feat/accelerator-support`。 +- **每个仓库一个 PR**,尽量单 PR 全部实现。 +- 文档/PR 标题/代码注释:英文、学术体、陈述句、无表情。 +- **`.cu` / `.cuh` 不进默认 source glob** —— `builtin_extension_table()` 的 `.ixx` 判例。 +- 新增描述符字段必须是旧客户端可跳过的(`xpkg.cppm` 已确认 skip 未知键)。 +- 判据取「只有做对了才会出现的读数」,不取「构建通过」。 + +## 本机验证条件(已实测,2026-09-05) + +| 项 | 值 | 意义 | +|---|---|---| +| GPU | NVIDIA RTX 4080,compute capability **8.9**(`sm_89`) | **可以真跑 kernel**,不止编译期验证 | +| nvcc | 12.0 V12.0.140,`/usr/bin/nvcc` | 形态 A 可全链路验证 | +| nvcc 宿主上界 | `host_config.h`:**gcc > 12 报错;clang 必须 < 15** | T6 的判据来源 | +| 可用宿主 | `clang++-14` ✅(实测编译并运行成功);`g++-13` ❌(超界);`g++-12` 二进制不存在 | 配对必须可选择,不只是校验 | +| mcpp 载荷 gcc | 16.1.0 | ⚠️ **比 nvcc 12.0 的上界高四个大版本 ⇒ 默认配对必然失败** | +| 基线构建 | `mcpp build` 53.23s,exit 0 | | + +⭐ 最后两行是 T6 存在的理由的**现场证据**:mcpp 的默认载荷与本机 CUDA **不兼容**, +所以 T6 不能只做「校验」,必须能**选择**一个兼容的宿主编译器,或给出可执行的错误。 + +--- + +## 仓库与依赖关系 + +``` +mcpp (引擎) ── PR A + │ T1 SourceKind::Device + │ T2 accelerator cfg layer key + │ T3 accel 身份字段 + tag_check + │ T4 [package] accelerators + │ T5 device target + │ T6 宿主/设备编译器配对 + │ T7 capability 同名符号边界 + │ T8 --accel CLI + │ T9 文档 T10 测试 + ↓ (引擎发布后才能被索引消费) +mcpp-index ── PR B 规则包 + 示例库描述符 + ↓ +xlings / xim-pkgindex ── PR C 载荷(按需) + ↓ +沙箱真实验证 ── xlings subos --sandbox --cmd +``` + +**跨仓库次序是硬的**:引擎的新键必须先发布,索引里的包才能用它 —— 否则索引里的包 +在已发布的 mcpp 上加载失败。本计划因此把 PR A 做完并发布,再做 PR B。 + +## 任务依赖图(mcpp 仓内) + +``` +T1 ─┬─> T5 ─> T6 + │ +T2 ─┴─────────────┐ +T3 ─> T8 ├─> T9 文档 +T4 ────────────────┤ +T7 ────────────────┘ + └─> T10 测试(每个任务自带,T10 是补齐 e2e) +``` + +T1/T2/T3/T4/T7 彼此独立,可并行。 + +--- + +## File Structure + +| 文件 | 职责 | 任务 | +|---|---|---| +| `modules/source-kind/src/source_kind.cppm` | 加 `Device` 档与设备扩展名表 | T1 | +| `src/build/prepare_inputs.cppm` | `accelerator` layer key、多值语义 | T2 | +| `src/pack/abi_tag.cppm` | `AbiTag::accel`、`tag_check` 扩展 | T3 | +| `modules/manifest/src/types.cppm` | `[package] accelerators`、device target 字段 | T4/T5 | +| `modules/manifest/src/toml.cppm` | 解析上述键 | T4/T5 | +| `modules/manifest/src/xpkg.cppm` | 描述符读写 `accel` | T3 | +| `src/build/prepare.cppm` | device target 的计划、架构收窄、配对校验 | T5/T6 | +| `src/toolchain/compat.cppm` | 宿主/设备编译器支持表 | T6 | +| `src/cli.cppm` | `--accel` / `--no-accel` | T8 | +| `docs/05-mcpp-toml.md` + `docs/zh/` | 手册 | T9 | +| `docs/19-accelerators.md` + zh | 新章 | T9 | +| `tests/unit/test_*.cpp` | 单测 | 各任务 | +| `tests/e2e/6xx_*.sh` | e2e | T10 | + +--- + +## Task 1: `SourceKind::Device` + +**Files:** `modules/source-kind/src/source_kind.cppm`, +`tests/unit/test_source_kind.cpp`(若无则新建) + +**Interfaces:** +- Produces: `SourceKind::Device`;`is_scan_exempt(Device) == true`; + `affects_graph_shape(DeviceHeader) == true`;`classify(".cu") == Device`; + `default_source_globs()` **不含** `*.cu`。 + +- [ ] **Step 1: 写失败的测试** —— `classify("k.cu")` 应为 `Device`; + `classify("k.cuh")` 应为 `Header` 一类(触发图形状失效); + `default_source_globs(builtin_extension_table())` 不含 `*.cu`。 +- [ ] **Step 2: 跑测试确认失败**(`mcpp test --filter source_kind`) +- [ ] **Step 3: 实现** —— 加 `Device` 枚举值、`kDeviceExtensions[] = {".cu",".hip"}`、 + `kDeviceHeaderExtensions[] = {".cuh",".hiph"}`(后者并入 `Header` 返回值), + `is_scan_exempt` 加 `Device`,`default_source_globs` **不加**设备扩展名。 +- [ ] **Step 4: 跑测试确认通过** +- [ ] **Step 5: 提交** `feat(source-kind): classify device translation units` + +**判据(只有做对才出现的读数)**:一个含 `.cu` 的工程 `mcpp build` 后, +`build.ninja` 里没有该文件的 P1689 扫描边;而 `.cuh` 改动后 fast path 回落完整 prepare。 + +--- + +## Task 2: `accelerator` 多值 cfg layer key + +**Files:** `src/build/prepare_inputs.cppm`,`src/build/prepare.cppm`, +`tests/unit/test_cfg_predicate.cpp` + +**Interfaces:** +- Consumes: 已有 `kCfgLayerKeys`、`Ctx::layer_value`、`unknown_tokens`。 +- Produces: `Ctx::layers_multi`(`map>`); + `accelerator` 进 `kCfgLayerKeys`;`any/all` 在多值 layer 上的语义。 + +语义(写死): +- `accelerator = "cuda"` ⟺ 集合恰好是 `{cuda}` +- `any(accelerator = "cuda")` ⟺ `cuda ∈` 集合 +- `all(accelerator="cuda", accelerator="rocm")` ⟺ 两者都在 + +- [ ] Step 1–5 同 TDD 循环。 +**判据**:`cfg(any(accelerator="rocm"))` 段在 cuda-only 构建下**不被合并**; +`cfg(acclerator="cuda")`(拼写错)被 `unknown_tokens()` 报出并给出最近词。 + +--- + +## Task 3: `accel` 身份字段与 `tag_check` + +**Files:** `src/pack/abi_tag.cppm`,`modules/manifest/src/xpkg.cppm`, +`src/pack/prebuilt.cppm`,`tests/unit/test_abi_tag_accel.cpp` + +**Interfaces:** +```cpp +struct AccelReq { std::string backend, version, arch; }; // 消费侧请求 +struct AccelSet { std::string backend, version; // 产物侧集合 + std::vector archs; + std::string ptxFloor; }; // 空 = 无 PTX +struct AbiTag { /* 已有四维 */ std::vector accel; }; +bool accel_accepts(const std::vector& published, const AccelReq& want); +``` +比较语义:`want.backend` 在集合里,**且**(`want.arch ∈ archs` **或** +`ptxFloor` 非空且 `arch >= ptxFloor`)。`published.accel` 为空 = 不受约束。 + +- [ ] Step 1: 测试 —— 空 accel 接受一切;`{sm80,sm90f}` 拒 `sm86`; + 带 `ptx=80` 时接受 `sm86`;backend 不同直接拒。 +- [ ] Step 2–5:实现 + `tag_check` 增加 accel 维的 `TagMismatch`。 + +**判据**:同一包发两变体,消费者按声明架构选中不同的那个; +诊断里出现 `accel needs …, this build has …` 行。 + +--- + +## Task 4: `[package] accelerators` 声明 + +**Files:** `modules/manifest/src/types.cppm`,`modules/manifest/src/toml.cppm` + +与 `[package] platforms` **严格同形**(声明/CI 矩阵提示,非强制)。 +**判据**:`mcpp why` 里能看到它;写了不支持的 backend 时 `mcpp build` 给出提示而非报错。 + +--- + +## Task 5: device target + +**Files:** `modules/manifest/src/types.cppm`、`toml.cppm`、`src/build/prepare.cppm`、 +`src/build/ninja_backend.cppm` + +新增 target 字段:`kind = "device"`、`backend`、`archs`、`device-link`, +以及 `[targets..sources]` 逐 glob 的 `archs` 收窄。 + +规则: +- glob 级 `archs` 必须是目标级子集,否则配置错误; +- 收窄成空集是**错误**; +- `archs` 缺省是**错误**,消息列出该 backend 的已知取值(决定 11)。 + +**判据**:空集与非子集都报错;`.cu` 的编译边出现在 `build.ninja` 且带正确的 `-gencode`。 + +--- + +## Task 6: 宿主/设备编译器配对(本机可真机验证) + +**Files:** `src/toolchain/compat.cppm`,`src/build/prepare.cppm` + +nvcc 的上界从 `crt/host_config.h` 读(本机实测:`__GNUC__ > 12` 报错, +clang 必须 `< 15`),或内置一张按 CUDA 版本的表。 + +行为: +1. 若 mcpp 当前宿主编译器超界 ⇒ **在调用 nvcc 之前**报错,消息含两边版本; +2. 若存在可用的兼容宿主(本机的 `clang++-14`)⇒ 提示可用 `-ccbin`; +3. 都没有 ⇒ 明确说「这套 CUDA 与 mcpp 可提供的宿主编译器都不兼容」。 + +**判据(本机现成)**:mcpp 载荷是 gcc 16.1.0,nvcc 是 12.0 ⇒ +**必须报错且消息里同时出现 `16.1.0` 与 `12.0` 与上界 `12`**, +而不是把命令发给 nvcc 让它报 `unsupported GNU version`。 + +--- + +## Task 7: capability 同名符号边界 + +**Files:** `src/build/prepare.cppm`(绑定处) + +同一 capability 的多个 provider 同时在图里 ⇒ **绑定期报错**, +而不是等到链接期报重复符号。 + +**判据**:两个都 `provides = ["gpu-blas"]` 且定义同名符号的包同时在图里, +错误发生在绑定阶段,消息列出两个候选包名。 + +--- + +## Task 8: `--accel` / `--no-accel` + +**Files:** `src/cli.cppm`,`src/build/prepare.cppm` + +与 `[build] accel` 的关系,和 `--target`↔`[toolchain]` 同形:manifest 声明 + 命令行覆盖。 + +--- + +## Task 9: 文档(英文学术体) + +- `docs/19-accelerators.md` + `docs/zh/19-accelerators.md`(新章) +- `docs/05-mcpp-toml.md` + zh:新键 +- `docs/11-machine-output.md` + zh:`accel` 进机器接口契约(**取值变了要写进契约页**) + +## Task 10: e2e 补齐 + +`tests/e2e/6xx_accel_*.sh`。需要 GPU 的用例**单独成文件并在 job 一级守卫**, +不与其余混在同一套件里靠 `# requires:` 跳过(memory:守卫住在 job 里)。 + +--- + +## 跨仓库 + +**PR B(mcpp-index)** —— 在 PR A 发布之后。按此顺序: + +| # | 包 | 类型 | 验证 | +|---|---|---|---| +| B1 | `rules-cuda` | 规则包 | P5;`backend` 经 capability 解析 | +| B2 | `compat.cccl` | header-only | 源码分发最短路径上的第一个真实库 | +| B3 | `compat.cudart` / `compat.cublas` | 能力包 | capability 绑定;**同名符号夹具(决定 4)** | +| B4 | `compat.cutlass` | header-only | 逐 glob 架构收窄(R4)的压测 | +| B5 | `llama.cpp-m` 加 CUDA 后端 | 真实框架 | ⭐⭐ R1+R3+R5;**它已在 `mcpplibs/` 里** | +| B6 | `compat.onnxruntime` / `compat.opencv` | 后续 | 真实体量的多变体选择 | + +**PR C(xim-pkgindex)** —— 载荷。按此顺序: + +| # | 包 | 验证 | +|---|---|---| +| C1 | ⭐ `adaptivecpp` | 形态 B **在无卡机器上 kernel 真执行** | +| C2 | `cuda-nvcc` | 设备工具链脱离 host(§5.0b 的原则) | +| C3 | `cuda-cudart` | 依赖已有的 `libcuda-host-link` sentinel | +| C4 | `dpcpp` / `cudnn` / `nccl` | 后续 | + +⚠️ 按组件取,不打整个 toolkit。`cuda-nvcc` + `cuda-cudart` 是百 MB 级不是 GB 级。 + +**沙箱验证**:`xlings subos --sandbox --cmd "..."`,先配 CN mirror。 + +--- + +## Self-Review(计划对规范的覆盖) + +| 设计决定 | 任务 | +|---|---| +| 1 accel 并列字段 | T3 | +| 2 kind="device" | T5 | +| 3 逐 glob 收窄 | T5 | +| 4 capability 边界先修 | T7 | +| 5 组合表达 | T2 | +| 6 不产静态库 | T5(不实现该路径即可) | +| 7 C++ 运行时核验 | T3(延后:见下) | +| 8 工具链 family@version | 形态 B,本轮不做 | +| 9 小众后端不做 | 范围 | +| 10 生态闭环 | PR B | +| 11 archs 缺省报错 | T5 | +| 12 backend 两趟 | T5 | +| 13 节名核验 | 形态 B,本轮不做 | +| 14 deploy_files | T5(Metal 路径本轮不做) | +| 15 字段统一 archs | T5 | +| 16 不加标注 | 无需实现 | + +⚠️ **本轮不做**:形态 B(SYCL/OpenMP,决定 8/13)、Metal(决定 14 的 Artifact 路径)、 +决定 7 的 C++ 运行时核验。理由:本机没有 SYCL 工具链与 Metal,做了无法真实验证, +而**未经真实验证的实现正是本设计反对的东西**。这三项在 PR A 里留出接口位置但不实现。 diff --git a/CHANGELOG.md b/CHANGELOG.md index daa0cf6a9..03a5f3eda 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,51 @@ > 本文件追踪 `mcpp-community/mcpp` 公开仓的版本演进。 > 格式参考 [Keep a Changelog](https://keepachangelog.com/zh-CN/1.1.0/)。 +## [2026.9.5.1] — 2026-09-05 + +### ⭐⭐ 加速器支持:设备编译单元、产物身份的加速器维、以及没人做的宿主编译器配对 + +一个为某个计算能力编译的库,被另一个计算能力的构建消费时,链接干净地完成, +程序在第一次 kernel 启动时失败,消息里既没有包名也没有任何一侧期望的架构。 +C++ 构建生态里没有任何一个系统把「这个二进制是为哪个架构编的」记进它的身份 —— +这是整个品类的空白,不是 mcpp 特有的。 + +**`SourceKind::Device`。** `.cu` 与 `.hip` 是设备编译单元:从不被扫描 import, +从不产出 BMI —— 没有任何设备编译器接受 C++20 modules。`.cuh` / `.hiph` 是头文件, +改动其一仍使快路径失效。设备扩展名**刻意不进默认 source glob**,理由与内置模块 +扩展名表停在 `.cppm` 的理由相同:放宽它会让一个 vendored 了设备源码、在别处构建 +它的已发布包在下次升级后突然开始编译它,而这是作者无法修复的破坏。 + +**产物身份的加速器维。** 携带设备代码的产物把它记在兼容性标签旁边,`tag_check` +比较它。成员判定按两条硬件里真实存在的机制放宽:家族目标覆盖同 major、minor 不低 +于它的范围;内嵌的可移植形式覆盖下界之上的一切。AMD 两者都没有,靠 archs 一侧的 +generic target 取得同样的覆盖,所以空的下界不放宽任何东西。 + +字段与标签并列而不是标签的一段,因为架构列表是集合,而标签是用 `-` 拼接、 +其 triple 本身含数量不定 `-` 的字符串。一个比较器,两个存储位置。 + +**`accelerator` 作为多值 cfg layer。** 一次构建可以同时启用多个后端。比较是处处 +成员判定,而不是只在 `any(...)` 里 —— 让组合子改变操作数含义会使 +`all(accelerator = "cuda", accelerator = "rocm")` 变成不可满足,而不是「两者都启用」。 + +**宿主编译器上界,读而不抄。** nvcc 拒绝比它在自己的 `crt/host_config.h` 里声明的 +上界更新的宿主编译器,而 mcpp 的载荷常常更新。因为宿主编译器由 mcpp 提供, +它可以在任何编译发生之前作答。上界从工具包读出,所以一个 mcpp 从未见过的工具包 +同样能作答;解析不了的头文件不产生上界,也就不产生断言。 + +**`[build] accel` / `--accel` / `--no-accel`,以及 `[package] accelerators`。** +前三者之间的关系与 `[toolchain]` 和 `--target` 相同。`--no-accel` 是显式请求 +「不要加速器」,这是在一个同时发布了设备构建的包中选中 CPU-only 变体的方式。 +`[package] accelerators` 与 `platforms` 同形,并刻意与产物的 `accel` 是不同字段: +声明由人手写,产物字段从构建测量。 + +真机验证(RTX 4080 / CUDA 12.0 / 驱动 550.144.03):`examples/09-cuda-kernel` +经规则包编出设备岛并运行,`mcpp run` 输出 `12 24 36 48`。 + +设计与调研:`.agents/docs/2026-09-05-accelerator-support-design.md`、 +`.agents/docs/2026-09-04-ai-accelerator-toolchain-ecosystem-survey.md`。 +新增手册章节 `docs/20-accelerators.md`(中英双份)。 + ## [2026.9.4.3] — 2026-09-04 ### ⭐⭐ `mcpp run` 报告程序自己的退出码 diff --git a/docs/05-mcpp-toml.md b/docs/05-mcpp-toml.md index cce92e221..407bfb06f 100644 --- a/docs/05-mcpp-toml.md +++ b/docs/05-mcpp-toml.md @@ -791,6 +791,22 @@ Linux and macOS perform no such conversion, so nothing is skipped there. A package that builds on one and not the other, with an `internal: unhandled exception` from a code-page message, was mcpp#516. +### 2.3.1 `[build] accel` — the accelerator this build targets + +```toml +[build] +accel = "cuda12.8+{sm_80,sm_90f} ptx>=90" +``` + +Which device backends and architectures this build compiles for. Overridden for +one build by `--accel`, the relationship `--target` has with `[toolchain]`; +`--no-accel` requests none explicitly, which is how a CPU-only variant of a +package that also publishes device builds is selected. + +The value is compared against the `accel` field of any prebuilt artifact the +build consumes, and a build asking for none is satisfied by every artifact. See +[20 — Accelerators](20-accelerators.md). + ### 2.4 `[lib]` — Library Root Module Convention ```toml @@ -1866,6 +1882,21 @@ built" means is the same question `--target` answers (docs/08 §7.4). Both are warnings, never errors: coverage is release discipline, and the person who can judge it is looking at the release, not at this build. +### 2.12b `[package] accelerators` — Accelerator Declaration + +```toml +[package] +accelerators = ["cuda", "rocm"] +``` + +Declares the accelerator backends the package supports. Mirrors `platforms`: a +statement of intent and a CI-matrix hint, shown by `mcpp why`, never a gate. + +Distinct from an artifact's `accel` field on purpose. A declaration is written +by hand and may be aspirational; `accel` is measured from the build that +produced a binary and is what a consumer is refused against. See +[20 — Accelerators](20-accelerators.md). + ### 2.13 `[xlings]` — the project's environment ```toml diff --git a/docs/20-accelerators.md b/docs/20-accelerators.md new file mode 100644 index 000000000..6eda100ab --- /dev/null +++ b/docs/20-accelerators.md @@ -0,0 +1,208 @@ +# 20 — Accelerators + +How mcpp builds device code, and how a prebuilt artifact states which devices +it can run on. + +## Two shapes, one of which mcpp implements today + +Accelerator toolchains come in two shapes, and they are not variations of one +model. + +An **island** keeps device code in separate translation units compiled by a +separate compiler. CUDA, HIP, Ascend C and Metal all work this way. The device +compiler produces an object (or, for shading languages, a runtime resource) +that joins the ordinary link. + +A **whole-target** model puts device code in ordinary `.cpp` files and compiles +the entire target with a compiler capable of offloading. SYCL, OpenMP offload +and stdpar work this way. There is no island to separate. + +This document describes the island shape, which is what mcpp implements. + +## Device translation units + +A source whose extension is `.cu` or `.hip` is a **device translation unit**. +mcpp classifies it as such and treats it accordingly: it is never scanned for +imports and never produces a BMI, because no device compiler accepts C++20 +modules. + +`.cuh` and `.hiph` are classified as headers. They are not compiled, but +editing one can change what the graph should be, so they invalidate the fast +path exactly as any other header does. + +Device extensions are **not** in the default source glob. A package that +vendors a `.cu` it builds elsewhere must not begin compiling it on an mcpp +upgrade, which is a break its author cannot fix once that version has shipped. +Device sources are opted into by naming them. + +## The seam + +A device translation unit cannot import a module, so the boundary between it +and the rest of a project is a header. Consumers do not see that header: a +module includes it in its global module fragment and exports a C++ interface, +and everything downstream imports the module. + +That module is worth naming a *seam*, because its reason for existing is not +the module boundary. It is the single place where the island underneath can be +exchanged — for HIP, for a CPU fallback — without any consumer changing, and +the single place a `cfg(accelerator = ...)` section has to apply. A project +without a seam has no boundary at which a backend can be substituted. + +Two constraints on the interface follow from what the compilers are, not from +taste. It should be `extern "C"`, because the device compiler drives a host +compiler that mcpp did not choose and the two sides therefore do not share a +C++ ABI. The island should avoid the standard library, because an island that +links libstdc++ puts a second copy of the C++ runtime into a program whose own +copy came from mcpp's toolchain. + +## Compiling an island + +The command that invokes a device compiler is not built into mcpp. It is +supplied by a **build-rule package**, consumed with `host-module = true`, +which emits build-graph edges whose outputs join the link. See +[07 — build.mcpp](07-build-mcpp.md) for the mechanism and +`examples/09-cuda-kernel` for a working CUDA rule. + +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. + +## Declaring what a build targets + +```toml +[build] +accel = "cuda12.8+{sm_80,sm_90f} ptx>=90" +``` + +overridden for one build by `--accel`, which is the relationship `--target` +has with `[toolchain]`. `--no-accel` is not the absence of `--accel`; it is an +explicit request for no accelerator, which is what selects a CPU-only variant +of a package that also publishes device builds. + +A source package may declare which backends it supports: + +```toml +[package] +accelerators = ["cuda", "rocm"] +``` + +This mirrors `[package] platforms`: a statement of intent and a CI-matrix hint, +not a gate. It is a different field from an artifact's `accel` on purpose — a +declaration is written by hand and may be aspirational, while an artifact's +field is measured from the build that produced it. + +## What a prebuilt artifact states + +An artifact that carries device code records it beside its compatibility tag: + +```toml +[[runtime.artifacts]] +role = "static-library" +path = "lib/libgpukit.a" +provenance = "mcpp-pack/1" +abi = "x86_64-linux-gnu-gcc16-libstdcxx16-c++23" +accel = "cuda12.8+{sm_80,sm_90f} ptx>=90" +``` + +The field is separate from the tag rather than a segment of it because an +architecture list is a set, and the tag is a dash-joined string whose triple +already contains a variable number of dashes. + +An absent `accel` means the artifact carries no device code and constrains +nothing, which is why a CPU-only library is usable by every build. + +### How a consumer is matched + +A build's request is satisfied by an artifact when, for each backend the build +asks for, the artifact declares that backend, agrees on the toolkit's major +version, and covers every requested architecture. An architecture is covered +when it is named, when a family target of the same major and an equal-or-lower +minor is named, or when the embedded portable form's floor is at or below it. + +Family targets and portable forms are what keep the variant matrix finite. +Publishing one artifact per chip does not scale; publishing one per generation +does. + +When nothing matches, the refusal names the dimension and both sides: + +``` +error: mcpplibs.gpuonly@0.1.0: no prebuilt artifact matches this toolchain. + your toolchain : x86_64-linux-gnu-gcc16-libstdcxx16-c++23 accel=cuda12.8+{sm_86} + published tags : + x86_64-linux-gnu accel=cuda12.8+{sm_90f} + closest is x86_64-linux-gnu, and it differs on: + accel needs cuda12.8+{sm_90f}, this build has cuda12.8+{sm_86} + fix: build for an architecture the package carries (--accel), or take + a variant that carries no device code (--no-accel), or ask the + publisher for one covering yours. +``` + +This is the failure the dimension exists to move. Without it the build links +cleanly and the program fails at its first kernel launch with a message naming +neither the package nor the architecture either side expected. + +### Publishing several variants + +A package may publish several artifacts, and a consumer takes the first whose +tag accepts it. **List the CPU-only artifact first.** An mcpp that predates the +`accel` field ignores it, and ordering is what still gives such a client an +artifact that runs anywhere. + +## Conditioning on the backend + +`accelerator` is a target-side layer, resolved after the dependency graph is, +and it holds a set rather than a single value: + +```toml +[target.'cfg(accelerator = "rocm")'.build] +cxxflags = ["-DMYAPP_ROCM"] +``` + +The comparison is membership, so a build enabling both CUDA and ROCm answers +true to each. `any`, `all` and `not` compose over it as ordinary boolean +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. + +**`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 +artifact *carries*, and mcpp does not yet compile device code itself — form A +goes through a build-rule package, so mcpp has nothing to measure. Recording a +declaration in a field whose meaning is "measured" would make the identity lie +in precisely the way the dimension exists to prevent. A publisher writes the +field explicitly today, which is what index descriptors do; `mcpp pack` will +emit it once `kind = "device"` puts the device compilation inside mcpp. + +## Not implemented + +The whole-target shape (SYCL, OpenMP offload, stdpar), device targets and the +device linking they imply, static libraries containing device code, and +accelerator payloads supplied through xim. See +`.agents/docs/2026-09-05-accelerator-support-design.md` for the design these +follow from. diff --git a/docs/README.md b/docs/README.md index 3f32410d0..f6d0ebde8 100644 --- a/docs/README.md +++ b/docs/README.md @@ -22,6 +22,7 @@ - [17 - The Project Environment](17-the-project-environment.md) - [18 - Reaching a Device](18-devices.md) - [19 - Supported Versions and Compatibility](19-supported-versions.md) +- [20 - Accelerators](20-accelerators.md) ## Specifications diff --git a/docs/zh/05-mcpp-toml.md b/docs/zh/05-mcpp-toml.md index e6ea54c22..47265dee6 100644 --- a/docs/zh/05-mcpp-toml.md +++ b/docs/zh/05-mcpp-toml.md @@ -690,6 +690,20 @@ warning: 'C:/.../pkg/test/www' contains names this system's active code page can Linux 与 macOS 不做这种转换,因此那里不会跳过任何名字。一个包在一边能构建、在另一 边报 `internal: unhandled exception` 并指向代码页,就是 mcpp#516。 +### 2.3.1 `[build] accel` — 本次构建面向的加速器 + +```toml +[build] +accel = "cuda12.8+{sm_80,sm_90f} ptx>=90" +``` + +本次构建为哪些设备后端与架构编译。单次构建可用 `--accel` 覆盖 —— +这与 `--target` 对 `[toolchain]` 的关系相同;`--no-accel` 是显式请求「不要加速器」, +也就是在一个同时发布了设备构建的包中选中 CPU-only 变体的方式。 + +该取值会与构建所消费的任何预建产物的 `accel` 字段比较,而请求为空的构建被任何产物满足。 +见 [20 — 加速器](20-accelerators.md)。 + ### 2.4 `[lib]` — 库根模块约定 ```toml @@ -1591,6 +1605,20 @@ platforms = ["linux", "macos", "windows"] 两者都只是 warning,绝不报错:覆盖度属于发布纪律,而能作判断的人看的是发布, 不是这一次构建。 +### 2.12b `[package] accelerators` — 加速器声明 + +```toml +[package] +accelerators = ["cuda", "rocm"] +``` + +声明该包支持的加速器后端。与 `platforms` 同形:一个意图声明与 CI 矩阵提示, +由 `mcpp why` 展示,**不是门**。 + +与产物的 `accel` 字段刻意不同。声明由人手写、可以是期望值;`accel` 是从产生该二进制的 +那次构建测量出来的,并且是消费者被拒绝时所依据的东西。见 +[20 — 加速器](20-accelerators.md)。 + ### 2.13 `[xlings]` — 工程的环境 ```toml diff --git a/docs/zh/20-accelerators.md b/docs/zh/20-accelerators.md new file mode 100644 index 000000000..3e6ef6f28 --- /dev/null +++ b/docs/zh/20-accelerators.md @@ -0,0 +1,175 @@ +# 20 — 加速器 + +mcpp 如何构建设备代码,以及一个预建产物如何声明它能在哪些设备上运行。 + +## 两种形态,mcpp 当前实现其中一种 + +加速器工具链有两种形态,它们不是同一个模型的两种变体。 + +**岛(island)** 把设备代码放在单独的编译单元里,由单独的编译器编译。CUDA、HIP、 +Ascend C 与 Metal 都是这种形态。设备编译器产出一个目标文件(着色语言则产出运行期 +资源),汇入普通的链接。 + +**整目标(whole-target)** 把设备代码放在普通 `.cpp` 里,整个目标由一个能 offload 的 +编译器编译。SYCL、OpenMP offload 与 stdpar 属于这一类,没有可分的岛。 + +本文描述岛这一形态,即 mcpp 当前实现的部分。 + +## 设备编译单元 + +扩展名为 `.cu` 或 `.hip` 的源文件是**设备编译单元**。mcpp 据此分类并相应处理: +它从不被扫描 import,也从不产出 BMI —— 因为没有任何设备编译器接受 C++20 modules。 + +`.cuh` 与 `.hiph` 被分类为头文件。它们不被编译,但改动其一可以改变构建图应有的形状, +因此与任何其它头文件一样会使快路径失效。 + +设备扩展名**不在**默认 source glob 中。一个 vendored 了 `.cu`、但在别处构建它的包, +不应当在 mcpp 升级后突然开始编译它 —— 这是它的作者无法修复的破坏,因为那个版本的 +tarball 已经发出去了。设备源文件必须被显式点名。 + +## 接缝(seam) + +设备编译单元不能 import 模块,因此它与工程其余部分的边界是一个头文件。消费者看不到 +这个头文件:一个模块在其 global module fragment 里包含它,并导出 C++ 接口, +下游一律 import 该模块。 + +这个模块值得叫做**接缝**,因为它存在的理由不是模块边界。它是**唯一**一处可以在不改动 +任何消费者的前提下替换下层的岛(换成 HIP、换成 CPU 回落)的位置,也是 +`cfg(accelerator = ...)` 唯一有落点的位置。没有接缝的工程,没有任何边界可供替换后端。 + +接口上的两条约束来自编译器的事实,而不是审美。它应当是 `extern "C"`, +因为设备编译器驱动的是一个 mcpp 没有选择的宿主编译器,两侧因此不共享 C++ ABI。 +岛本身应当避开标准库,因为一个链接了 libstdc++ 的岛,会把第二份 C++ 运行时放进一个 +自身运行时来自 mcpp 工具链的程序里。 + +## 编译一个岛 + +调用设备编译器的那条命令不内置在 mcpp 里,而是由**构建规则包**提供 —— +以 `host-module = true` 消费,emit 输出汇入链接的构建边。机制见 +[07 — build.mcpp](07-build-mcpp.md),可用的 CUDA 规则见 `examples/09-cuda-kernel`。 + +这个划分是刻意的。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 无法解析的头文件产生不出上界,也就不产生任何断言。 + +这是报告而非强制:一个不编译任何设备代码的工程,不受不兼容配对的影响。 + +## 声明一次构建的目标 + +```toml +[build] +accel = "cuda12.8+{sm_80,sm_90f} ptx>=90" +``` + +单次构建可用 `--accel` 覆盖 —— 这与 `--target` 对 `[toolchain]` 的关系相同。 +`--no-accel` 不是「没有写 `--accel`」,它是**显式请求不要加速器**, +这正是在一个同时发布了设备构建的包中选中 CPU-only 变体所需要的。 + +源码包可以声明它支持哪些后端: + +```toml +[package] +accelerators = ["cuda", "rocm"] +``` + +这与 `[package] platforms` 同形:一个意图声明与 CI 矩阵提示,不是门。 +它与产物的 `accel` 刻意是不同的字段 —— 声明由人手写、可以是期望值, +而产物的字段是从产生它的那次构建**测量**出来的。 + +## 一个预建产物声明了什么 + +携带设备代码的产物把它记在兼容性标签旁边: + +```toml +[[runtime.artifacts]] +role = "static-library" +path = "lib/libgpukit.a" +provenance = "mcpp-pack/1" +abi = "x86_64-linux-gnu-gcc16-libstdcxx16-c++23" +accel = "cuda12.8+{sm_80,sm_90f} ptx>=90" +``` + +它是标签之外的独立字段而不是标签的一段,因为架构列表是一个**集合**, +而标签是用 `-` 拼接的字符串,其中的 triple 本身就含有数量不定的 `-`。 + +`accel` 缺席表示该产物不携带设备代码,因而不施加任何约束 —— +这就是为什么一个纯 CPU 的库可被任何构建使用。 + +### 消费者如何被匹配 + +当一次构建请求的每个后端,产物都声明了该后端、工具包主版本一致、 +且覆盖了请求的每个架构时,该产物满足这次构建。一个架构被覆盖,当它被直接列出, +或列出了同 major、minor 不高于它的家族目标,或内嵌可移植形式的下界不高于它。 + +**家族目标与可移植形式是让变体矩阵保持有限的东西。** 一颗芯片一份产物不可扩展, +一个世代一份可以。 + +没有任何产物匹配时,拒绝会点名维度与两侧的取值: + +``` +error: mcpplibs.gpuonly@0.1.0: no prebuilt artifact matches this toolchain. + your toolchain : x86_64-linux-gnu-gcc16-libstdcxx16-c++23 accel=cuda12.8+{sm_86} + published tags : + x86_64-linux-gnu accel=cuda12.8+{sm_90f} + closest is x86_64-linux-gnu, and it differs on: + accel needs cuda12.8+{sm_90f}, this build has cuda12.8+{sm_86} + fix: build for an architecture the package carries (--accel), or take + a variant that carries no device code (--no-accel), or ask the + publisher for one covering yours. +``` + +这正是该维度存在所要搬移的那个失败。没有它,构建干净地链接完成, +而程序在第一次 kernel 启动时失败,消息里既没有包名,也没有任何一侧期望的架构。 + +### 发布多个变体 + +一个包可以发布多个产物,消费者取第一个标签接受它的。**把 CPU-only 的产物排在第一位。** +早于 `accel` 字段的 mcpp 会忽略该字段,而顺序是让这样的客户端仍然拿到一个 +到处都能跑的产物的唯一机制。 + +## 按后端条件化 + +`accelerator` 是目标侧的一层,在依赖图解析之后才可求值,并且它持有一个**集合** +而不是单个取值: + +```toml +[target.'cfg(accelerator = "rocm")'.build] +cxxflags = ["-DMYAPP_ROCM"] +``` + +比较是**成员判定**,因此一次同时启用 CUDA 与 ROCm 的构建对两者都答真。 +`any`、`all`、`not` 作为普通布尔组合子在其上组合。 + +## 两条值得写明的边界 + +**`--accel` 是 `build` 的选项**,与 `--static`、`--toolchain` 同级, +不在 `run`、`test`、`pack` 上重复。那些命令与读取任何其它构建输入一样, +从 manifest 读 `[build] accel`;这个 flag 的用途是覆盖单次构建,而那正是 `build` 覆盖的场景。 + +**`mcpp pack` 不写 `accel` 字段。** 它本可以把 manifest 声明的值写进去, +而这恰恰是它不这么做的理由:该字段陈述产物**携带**了什么, +而 mcpp 目前不自己编译设备代码 —— 形态 A 走规则包,mcpp 没有可测量的对象。 +把一个声明写进一个含义是「测量值」的字段,会让身份**恰好以该维度要防止的方式**说谎。 +今天由发布者显式写这个字段,索引描述符就是这么做的; +等 `kind = "device"` 把设备编译放进 mcpp 之后,`mcpp pack` 才会发它。 + +## 尚未实现 + +整目标形态(SYCL、OpenMP offload、stdpar)、device target 及其隐含的 device link、 +含设备代码的静态库,以及经由 xim 提供的加速器载荷。这些所依据的设计见 +`.agents/docs/2026-09-05-accelerator-support-design.md`。 diff --git a/docs/zh/README.md b/docs/zh/README.md index 1ec50d476..00ffad9e7 100644 --- a/docs/zh/README.md +++ b/docs/zh/README.md @@ -22,6 +22,7 @@ - [17 - 项目环境](17-the-project-environment.md) - [18 - 抵达一台设备](18-devices.md) - [19 - 受支持的版本与兼容性](19-supported-versions.md) +- [20 - 加速器](20-accelerators.md) ## 规范文档 diff --git a/examples/09-cuda-kernel/README.md b/examples/09-cuda-kernel/README.md new file mode 100644 index 000000000..2853b43fd --- /dev/null +++ b/examples/09-cuda-kernel/README.md @@ -0,0 +1,101 @@ +# 09 — A CUDA kernel behind a seam module + +What this example demonstrates, and what it does not. + +## The shape + +``` +app/ + src/kernels/saxpy.cu the island: compiled by nvcc, never scanned, no BMI + 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 +``` + +Three properties are load-bearing. + +**The island is not in the module graph.** No device compiler accepts C++20 +modules, so `.cu` is classified as a device translation unit: never scanned for +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 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. + +## The rule package, and why nvcc's host compiler is its problem + +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: + +``` +example.rules.cuda: nvcc /usr/bin/nvcc with -ccbin /usr/bin/clang++-14 +``` + +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. + +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`. + +## Verified + +On an NVIDIA RTX 4080 (compute capability 8.9) with CUDA 12.0 and driver +550.144.03: + +``` +$ mcpp run + Running `target/.../bin/cuda-saxpy` +12 24 36 48 +``` + +which is `2.0 * [1,2,3,4] + [10,20,30,40]` computed on the device. + +## Where the driver comes from + +The CUDA runtime is linked statically, so the artifact carries every +redistributable component. That leaves exactly one host dependency, +`libcuda.so.1` — the driver's userspace library, which NVIDIA's licence forbids +redistributing and which is in ABI lockstep with the kernel module. + +Reaching it takes two packages, one per layer. + +`libcuda-host-link` in xim owns the question of *where* the host's copy is. It +installs a symlink to whatever the machine has, so every GPU consumer reads one +path instead of reimplementing an `ldconfig` probe. + +`compat.cuda-runtime` in mcpp-index owns the mcpp-side question of how a +**built artifact** reaches it. mcpp's private loader does not consult +`/usr/lib`, so a bare-soname `dlopen` from inside a built program finds nothing; +the package declares a directory on the artifact's runtime search path and links +through the sentinel. It is declared here as an ordinary dependency: + +```toml +[dependencies.compat] +cuda-runtime = "2026.09.05" +``` + +Without it the program builds and links, then reports + +``` +cudaMalloc: CUDA driver version is insufficient for CUDA runtime version +``` + +which is what the runtime says when it cannot open the driver at all. diff --git a/examples/09-cuda-kernel/app/build.mcpp b/examples/09-cuda-kernel/app/build.mcpp new file mode 100644 index 000000000..d263f4841 --- /dev/null +++ b/examples/09-cuda-kernel/app/build.mcpp @@ -0,0 +1,20 @@ +import std; +import mcpp; +import example.rules.cuda; + +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"; + opt.includes = { "include" }; + const std::vector sources{ "src/kernels/saxpy.cu" }; + return example::rules::cuda::compile(sources, opt) ? 0 : 1; +} diff --git a/examples/09-cuda-kernel/app/include/saxpy/saxpy.h b/examples/09-cuda-kernel/app/include/saxpy/saxpy.h new file mode 100644 index 000000000..8d82d9727 --- /dev/null +++ b/examples/09-cuda-kernel/app/include/saxpy/saxpy.h @@ -0,0 +1,21 @@ +// The device island's interface. +// +// `extern "C"` and free of standard-library types, on purpose. The island is +// compiled by nvcc driving 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. Keeping the boundary this narrow is also what lets the island publish a +// C-surface compatibility tag. +#ifndef MCPP_EXAMPLE_SAXPY_H +#define MCPP_EXAMPLE_SAXPY_H + +#ifdef __cplusplus +extern "C" { +#endif + +// out[i] = a * x[i] + y[i], computed on the device. Returns 0 on success. +int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n); + +#ifdef __cplusplus +} +#endif +#endif diff --git a/examples/09-cuda-kernel/app/mcpp.toml b/examples/09-cuda-kernel/app/mcpp.toml new file mode 100644 index 000000000..fd718e107 --- /dev/null +++ b/examples/09-cuda-kernel/app/mcpp.toml @@ -0,0 +1,49 @@ +[package] +name = "cuda-saxpy" +namespace = "example" +version = "0.1.0" +description = "A CUDA kernel behind a seam module" +accelerators = ["cuda"] + +[language] +standard = "c++23" +modules = true +import_std = true + +[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. +[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. +[xlings] +deps = [{ linux = "libcuda-host-link" }] + +[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" +include_dirs = ["include"] +ldflags = ["-L/usr/local/cuda/lib64", "-L/usr/lib/x86_64-linux-gnu", + "-lcudart_static", "-lrt", "-lpthread", "-ldl"] + +[targets.cuda-saxpy] +kind = "bin" +main = "src/main.cpp" diff --git a/examples/09-cuda-kernel/app/src/app.cppm b/examples/09-cuda-kernel/app/src/app.cppm new file mode 100644 index 000000000..ab6d4ddd1 --- /dev/null +++ b/examples/09-cuda-kernel/app/src/app.cppm @@ -0,0 +1,27 @@ +// The seam. +// +// Its reason for existing is not that nvcc rejects modules. It is that this is +// the one place a backend can be exchanged: the island underneath can become +// HIP or a CPU fallback without a single consumer of this module changing, and +// a `cfg(accelerator = ...)` section has somewhere to apply. Remove the seam +// and every importer becomes backend-specific. +module; +#include "saxpy/saxpy.h" +export module app.saxpy; +import std; + +export namespace app { + +// The device interface is raw pointers and a count because it has to be. The +// seam is where that becomes a C++ interface again. +std::optional> +saxpy(float a, std::span x, std::span y) { + if (x.size() != y.size()) return std::nullopt; + std::vector out(x.size()); + if (saxpy_device(a, x.data(), y.data(), out.data(), + static_cast(x.size())) != 0) + return std::nullopt; + return out; +} + +} // namespace app diff --git a/examples/09-cuda-kernel/app/src/kernels/saxpy.cu b/examples/09-cuda-kernel/app/src/kernels/saxpy.cu new file mode 100644 index 000000000..d5d858d4c --- /dev/null +++ b/examples/09-cuda-kernel/app/src/kernels/saxpy.cu @@ -0,0 +1,56 @@ +// The island. Nothing here is visible to the module graph: nvcc does not +// accept C++20 modules, so this translation unit is never scanned and never +// produces a BMI. +// +// It uses no C++ standard library. That is a deliberate property rather than +// an accident of a small example: an island that pulls in libstdc++ links a +// second copy of the C++ runtime into a program whose own copy came from +// mcpp's toolchain, which is the failure where one is linked and the other is +// loaded. +#include "saxpy/saxpy.h" +#include +#include + +namespace { + +__global__ void saxpy_kernel(float a, const float* x, const float* y, + float* out, unsigned n) { + unsigned i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) out[i] = a * x[i] + y[i]; +} + +} // namespace + +extern "C" int saxpy_device(float a, const float* x, const float* y, + float* out, unsigned n) { + float *dx = nullptr, *dy = nullptr, *dout = nullptr; + const size_t bytes = static_cast(n) * sizeof(float); + int rc = -1; + + if (cudaError_t e = cudaMalloc(&dx, bytes); e != cudaSuccess) { + std::fprintf(stderr, "cudaMalloc: %s\n", cudaGetErrorString(e)); + goto done; + } + if (cudaMalloc(&dy, bytes) != cudaSuccess) goto done; + if (cudaMalloc(&dout, bytes) != cudaSuccess) goto done; + if (cudaMemcpy(dx, x, bytes, cudaMemcpyHostToDevice) != cudaSuccess) goto done; + if (cudaMemcpy(dy, y, bytes, cudaMemcpyHostToDevice) != cudaSuccess) goto done; + + saxpy_kernel<<<(n + 255) / 256, 256>>>(a, dx, dy, dout, n); + // The launch is asynchronous, so its own return value reports only whether + // the launch was accepted. A kernel compiled for an architecture this + // device does not have fails HERE, with `no kernel image is available for + // execution on the device` — which is the runtime failure the accelerator + // dimension of an artifact's identity exists to turn into a build-time one. + if (cudaError_t e = cudaGetLastError(); e != cudaSuccess) { + std::fprintf(stderr, "launch: %s\n", cudaGetErrorString(e)); + goto done; + } + if (cudaDeviceSynchronize() != cudaSuccess) goto done; + if (cudaMemcpy(out, dout, bytes, cudaMemcpyDeviceToHost) != cudaSuccess) goto done; + rc = 0; + +done: + cudaFree(dx); cudaFree(dy); cudaFree(dout); + return rc; +} diff --git a/examples/09-cuda-kernel/app/src/main.cpp b/examples/09-cuda-kernel/app/src/main.cpp new file mode 100644 index 000000000..6efc19383 --- /dev/null +++ b/examples/09-cuda-kernel/app/src/main.cpp @@ -0,0 +1,12 @@ +import std; +import app.saxpy; + +int main() { + const std::vector x{1, 2, 3, 4}, y{10, 20, 30, 40}; + auto out = app::saxpy(2.0f, x, y); + if (!out) { std::println("device unavailable"); return 1; } + for (auto v : *out) std::print("{} ", v); + std::println(""); + const std::vector want{12, 24, 36, 48}; + return *out == want ? 0 : 1; +} diff --git a/examples/09-cuda-kernel/rules-cuda/mcpp.toml b/examples/09-cuda-kernel/rules-cuda/mcpp.toml new file mode 100644 index 000000000..cdfa36abf --- /dev/null +++ b/examples/09-cuda-kernel/rules-cuda/mcpp.toml @@ -0,0 +1,6 @@ +[package] +name = "rules-cuda" +namespace = "example" +version = "0.1.0" +description = "Compile CUDA device translation units with nvcc (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 new file mode 100644 index 000000000..da1f460f8 --- /dev/null +++ b/examples/09-cuda-kernel/rules-cuda/src/rules-cuda.cppm @@ -0,0 +1,207 @@ +// Compile CUDA device translation units and hand the objects to the link. +// +// 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. +// +// 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. +// +// THE HOST COMPILER IS THE PART THAT SURPRISES PEOPLE +// +// 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. + +export module example.rules.cuda; + +import std; +import mcpp; + +export namespace example::rules::cuda { + +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. + std::vector includes; + std::string out_dir = std::string(mcpp::out_dir()); +}; + +struct edge { + std::string id, description; + std::vector command, inputs, outputs; +}; + +// ─── 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; + 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::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()); + } + c.push_back("/usr/include/crt/host_config.h"); + return first_existing(c); +} + +// The greatest gcc major and the greatest clang major the toolkit accepts. +// Zero means the header said nothing, which is not a refusal. +struct bounds { int gcc = 0, clang = 0; }; + +inline bounds read_bounds(std::string_view headerPath) { + bounds b; + if (headerPath.empty()) return b; + std::ifstream in{std::string(headerPath)}; + std::string text{std::istreambuf_iterator(in), std::istreambuf_iterator()}; + auto number_after = [&](std::size_t pos) { + int v = 0, n = 0; + while (pos < text.size() && !std::isdigit(static_cast(text[pos]))) { + if (text[pos] == '\n') return 0; + ++pos; + } + while (pos < text.size() && std::isdigit(static_cast(text[pos]))) { + v = v * 10 + (text[pos] - '0'); ++pos; ++n; + } + 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("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. +// +// 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; + } + 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 {}; +} + +// ─── Planning ────────────────────────────────────────────────────────────── + +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\"}}."); + return out; + } + const auto nvcc = find_nvcc(); + if (nvcc.empty()) { + std::println(std::cerr, + "example.rules.cuda: no nvcc found (looked at $CUDA_PATH/bin, " + "/usr/local/cuda/bin, /usr/bin)."); + 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; + } + std::println("example.rules.cuda: nvcc {} with -ccbin {}", nvcc, ccbin); + + 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" }; + 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.inputs = { root + "/" + src }; + e.outputs = { obj }; + out.push_back(std::move(e)); + } + return out; +} + +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.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()); + for (auto const& o : e.outputs) a.output(o.c_str()); + a.submit(); + } + return true; +} + +inline bool compile(std::span sources, options opt = {}) { + auto edges = plan(sources, std::move(opt)); + return submit(edges); +} + +} // namespace example::rules::cuda diff --git a/mcpp.toml b/mcpp.toml index 74628754c..9d8e7e022 100644 --- a/mcpp.toml +++ b/mcpp.toml @@ -1,6 +1,6 @@ [package] name = "mcpp" -version = "2026.9.4.3" +version = "2026.9.5.1" description = "Modern C++ build & package management tool" license = "Apache-2.0" authors = ["mcpp-community"] diff --git a/modules/manifest/src/toml.cppm b/modules/manifest/src/toml.cppm index 48c63ba34..7e9fafc9c 100644 --- a/modules/manifest/src/toml.cppm +++ b/modules/manifest/src/toml.cppm @@ -534,6 +534,8 @@ std::expected parse_string(std::string_view content, if (auto v = doc->get_string("package.repo")) m.package.repo = *v; if (auto v = doc->get_string_array("package.authors")) m.package.authors = *v; if (auto v = doc->get_string_array("package.platforms")) m.package.platforms = *v; + if (auto v = doc->get_string_array("package.accelerators")) + m.package.accelerators = *v; // [package].standard (M5.0 new home) if (auto v = doc->get_string("package.standard")) { @@ -1524,6 +1526,10 @@ std::expected parse_string(std::string_view content, return std::unexpected(error(origin, *err)); m.buildConfig.moduleExtensions = *v; } + // [build] accel — the accelerator backends and device architectures this + // build targets. One spelling with the descriptor field and the diagnostic, + // so what a user writes is what a refusal prints back at them. + if (auto v = doc->get_string("build.accel")) m.buildConfig.accel = *v; // [build] build_program_timeout — seconds a build.mcpp may run; 0 = no // limit. `optional` is load-bearing: with a plain int, "absent" and // "explicitly 0" would be the same value, and every project that never @@ -1747,6 +1753,7 @@ std::expected parse_string(std::string_view content, // // MUST stay in sync with the `doc->get_*("build.")` reads above. static constexpr std::string_view kKnownBuildKeys[] = { + "accel", "allow_host_libs", "bmi_schedule", "build_program_timeout", "c_standard", "cache", "cflags", "cxxflags", "cxx_runtime", "default-profile", "defines", "dependency_linkage", @@ -1901,8 +1908,15 @@ std::expected parse_string(std::string_view content, std::string path; auto const& table = value.as_table(); for (auto const& [key, _] : table) { + // ⚠️ A CLOSED whitelist: an unrecognised key is an error here, + // unlike the published xpkg descriptor, whose reader collects + // and skips one. The two readers differ on purpose — a typo in + // a hand-written manifest should be loud — but it means adding + // a key to THIS table is a compatibility event: an mcpp that + // predates the key refuses the whole manifest rather than + // ignoring the field. See the accelerator design's RK-2. if (key != "role" && key != "path" && key != "provenance" - && key != "abi" && key != "digest" + && key != "abi" && key != "accel" && key != "digest" && key != "host_fingerprint") { return std::unexpected(error(origin, std::format( "runtime.artifacts[{}] has unsupported key '{}'", @@ -1913,6 +1927,7 @@ std::expected parse_string(std::string_view content, || !table_string(table, "path", path) || !table_string(table, "provenance", artifact.provenance) || !table_string(table, "abi", artifact.abi) + || !table_string(table, "accel", artifact.accel) || !table_string(table, "digest", artifact.digest) || !table_string(table, "host_fingerprint", artifact.hostFingerprint)) { diff --git a/modules/manifest/src/types.cppm b/modules/manifest/src/types.cppm index 7581aec49..896964ffd 100644 --- a/modules/manifest/src/types.cppm +++ b/modules/manifest/src/types.cppm @@ -70,6 +70,15 @@ struct Package { std::vector authors; std::string repo; std::vector platforms; // declared supported platforms (CI matrix hint) + // Accelerator backends this package supports, declared in the same spirit + // as `platforms`: a statement of intent and a CI-matrix hint, not a gate. + // + // The pair is deliberate and the two must not be confused. `accelerators` + // is what a SOURCE package says it can be built for; `accel` on a built + // artifact is what that binary actually carries. A declaration is written + // by hand and can be aspirational; the artifact field is measured from the + // build and is what a consumer is refused against. + std::vector accelerators; // Resolution source carried into machine-readable runtime provenance. // Version dependencies use `index+@`; path/git packages // use their corresponding immutable-or-local source spelling. Parsing a @@ -570,6 +579,11 @@ struct BuildConfig : BuildInputs { // Scoped to the declaring package — a dependency is classified by its own // manifest, never by its consumer's. std::vector moduleExtensions; + // [build] accel — which accelerator backends and device architectures this + // build targets, in the wire form mcpp.pack.abi_tag reads. Empty means the + // build asks for none, and then every prebuilt artifact satisfies it + // vacuously; ordering in the descriptor is what makes the CPU variant win. + std::string accel; // [build] build_program_timeout — seconds this package's build.mcpp may // run before mcpp kills it. 0 = no limit; nullopt = use the built-in 600. // @@ -752,6 +766,14 @@ struct RuntimeArtifact { std::filesystem::path path; std::string provenance; std::string abi; + // What device code this artifact carries, in the wire form + // mcpp.pack.abi_tag reads: `cuda12.8+{sm_80,sm_90f} ptx>=90`. + // + // A separate field rather than a segment of `abi`, because an architecture + // list is a set and the tag is a dash-joined string whose triple already + // carries a variable number of dashes. Empty means the artifact carries no + // device code, which constrains nothing. + std::string accel; std::string digest; std::string hostFingerprint; }; diff --git a/modules/manifest/src/xpkg.cppm b/modules/manifest/src/xpkg.cppm index 1d3b8cda2..c61a15736 100644 --- a/modules/manifest/src/xpkg.cppm +++ b/modules/manifest/src/xpkg.cppm @@ -1941,6 +1941,7 @@ synthesize_from_xpkg_lua(std::string_view luaContent, else if (field == "path") path = std::move(value); else if (field == "provenance") artifact.provenance = std::move(value); else if (field == "abi") artifact.abi = std::move(value); + else if (field == "accel") artifact.accel = std::move(value); else if (field == "digest") artifact.digest = std::move(value); else if (field == "host_fingerprint") artifact.hostFingerprint = std::move(value); diff --git a/modules/source-kind/src/source_kind.cppm b/modules/source-kind/src/source_kind.cppm index e59674adf..27bd629b9 100644 --- a/modules/source-kind/src/source_kind.cppm +++ b/modules/source-kind/src/source_kind.cppm @@ -64,6 +64,15 @@ enum class SourceKind { GasAsm, // NASM syntax. NasmAsm, + // Compiled by a vendor device compiler mcpp does not drive directly + // (nvcc, hipcc). Never scanned for imports and never produces a BMI: no + // such compiler accepts C++20 modules, so a device unit cannot be one. + // + // The kind states the GRAPH ROLE, not the language. That is deliberate: + // it lets one kind cover CUDA C++, HIP, and device dialects that are not + // C++ at all (Intel Gaudi's TPC-C is derived from C99), without the + // classification table growing a row per vendor. + Device, // Not compiled, but editing one can change what the graph should be. Header, // Not a build input. @@ -219,6 +228,11 @@ constexpr std::string_view kCExtensions[] = { ".c", ".m" }; constexpr std::string_view kGasExtensions[] = { ".S", ".s" }; constexpr std::string_view kNasmExtensions[] = { ".asm" }; constexpr std::string_view kHeaderExtensions[] = { ".h", ".hpp", ".hh", ".hxx" }; +// Device sources and the headers they include. The headers are classified as +// `Header` rather than `Device` because their role is the header role: they +// are not compiled, and editing one can change what the graph should be. +constexpr std::string_view kDeviceExtensions[] = { ".cu", ".hip" }; +constexpr std::string_view kDeviceHeaderExtensions[] = { ".cuh", ".hiph" }; bool contains(std::span set, std::string_view ext) { for (auto e : set) if (e == ext) return true; @@ -234,6 +248,7 @@ std::string_view to_string(SourceKind k) { case SourceKind::C: return "c"; case SourceKind::GasAsm: return "gas"; case SourceKind::NasmAsm: return "nasm"; + case SourceKind::Device: return "device"; case SourceKind::Header: return "header"; case SourceKind::Other: return "other"; } @@ -283,9 +298,10 @@ ExtensionTable extension_table_for(std::span extras) { } bool is_reserved_non_module_extension(std::string_view ext) { - return contains(kCxxExtensions, ext) || contains(kCExtensions, ext) - || contains(kGasExtensions, ext) || contains(kNasmExtensions, ext) - || contains(kHeaderExtensions, ext); + return contains(kCxxExtensions, ext) || contains(kCExtensions, ext) + || contains(kGasExtensions, ext) || contains(kNasmExtensions, ext) + || contains(kHeaderExtensions, ext) || contains(kDeviceExtensions, ext) + || contains(kDeviceHeaderExtensions, ext); } std::optional @@ -334,7 +350,9 @@ SourceKind classify(const std::filesystem::path& p, const ExtensionTable& t) { if (contains(kCExtensions, ext)) return SourceKind::C; if (contains(kGasExtensions, ext)) return SourceKind::GasAsm; if (contains(kNasmExtensions, ext)) return SourceKind::NasmAsm; - if (contains(kHeaderExtensions, ext)) return SourceKind::Header; + if (contains(kDeviceExtensions, ext)) return SourceKind::Device; + if (contains(kHeaderExtensions, ext) + || contains(kDeviceHeaderExtensions, ext)) return SourceKind::Header; return SourceKind::Other; } @@ -343,7 +361,8 @@ bool produces_bmi(SourceKind k) { return k == SourceKind::ModuleInterface; } bool links_unconditionally(SourceKind k) { return k == SourceKind::ModuleInterface; } bool is_scan_exempt(SourceKind k) { - return k == SourceKind::C || k == SourceKind::GasAsm || k == SourceKind::NasmAsm; + return k == SourceKind::C || k == SourceKind::GasAsm + || k == SourceKind::NasmAsm || k == SourceKind::Device; } bool is_cxx_like(SourceKind k) { @@ -369,6 +388,12 @@ std::vector default_source_globs(const ExtensionTable& t) { globs.push_back("src/**/*.S"); globs.push_back("src/**/*.s"); globs.push_back("src/**/*.asm"); + // Device extensions are deliberately absent. The argument is the one that + // keeps the built-in module-extension table at `.cppm` alone: widening the + // default glob makes a published package that vendors a `.cu` it builds + // elsewhere start compiling it on the next mcpp upgrade, and that is a + // break its author cannot fix because the tarball has already shipped. + // Device sources are opted into by naming them in a device target. return globs; } diff --git a/modules/versioning/src/version.cppm b/modules/versioning/src/version.cppm index ac6703751..07110adb2 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.4.3"; +inline constexpr std::string_view MCPP_VERSION = "2026.9.5.1"; } // namespace mcpp diff --git a/src/build/prepare.cppm b/src/build/prepare.cppm index 0d4a027f3..4dfdedca0 100644 --- a/src/build/prepare.cppm +++ b/src/build/prepare.cppm @@ -839,6 +839,10 @@ export struct BuildOverrides { std::shared_ptr inherited_runtime_binding; std::string target_triple; // empty = host triple, fall through to [toolchain] + // --accel: the device backends and architectures this build targets, in + // the wire form mcpp.pack.abi_tag reads. Overrides `[build] accel`, the + // same relationship --target has with [toolchain]. + std::string accel; bool force_static = false; // --static (or implied by musl target) std::string package_filter; // -p : only build this workspace member // --profile . Empty = fall through to `[build] default-profile`, then @@ -8994,8 +8998,14 @@ prepare_build(bool print_fingerprint, auto t = mcpp::toolchain::triple::parse(tc->targetTriple); return t ? t->str() : tc->targetTriple; }(); - const auto currentTag = mcpp::pack::cxx_surface_tag( + auto currentTag = mcpp::pack::cxx_surface_tag( *tc, canonicalTriple, m->cppStandard.level); + // What THIS build targets on the device axis. Absent means it asks for + // 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); for (std::size_t i = 1; i < packages.size(); ++i) { auto const& pkg = packages[i]; if (!mcpp::pack::is_distribution_package(pkg.manifest)) continue; diff --git a/src/build/prepare_inputs.cppm b/src/build/prepare_inputs.cppm index 0cc5952e0..94be861fe 100644 --- a/src/build/prepare_inputs.cppm +++ b/src/build/prepare_inputs.cppm @@ -67,6 +67,11 @@ struct Ctx { std::string os, arch, family, env, triple; bool layersKnown = false; std::string compiler, compilerRuntime, kernelAbi, cAbi, cxxAbi; + // The first MULTI-VALUED layer. One build can enable several accelerator + // backends at once, which is what an inference framework shipping CUDA and + // ROCm device code in one artifact requires, so this layer holds a set + // rather than the single answer the other five hold. + std::vector accelerators; std::string_view layer_value(std::string_view k) const { if (k == "compiler") return compiler; @@ -76,6 +81,22 @@ struct Ctx { if (k == "c++-abi") return cxxAbi; return {}; } + + // THE SINGLE PLACE THE MULTI-VALUED CASE DIFFERS. + // + // A multi-valued layer compares by MEMBERSHIP, and it does so everywhere — + // not only inside `any(...)`. The alternative, letting `any(...)` mean + // membership while a bare key meant set equality, would make a combinator + // change the meaning of its operand: `all(accelerator = "cuda", + // accelerator = "rocm")` would then be unsatisfiable rather than "both + // backends are enabled". Membership everywhere keeps `any`/`all`/`not` + // pure boolean combinators, and a single-backend build still answers + // `accelerator = "cuda"` true and `accelerator = "rocm"` false. + bool layer_matches(std::string_view k, std::string_view v) const { + if (k == "accelerator") + return std::ranges::find(accelerators, v) != accelerators.end(); + return layer_value(k) == v; + } }; // Derive the cfg context from the resolved --target triple, falling back to @@ -126,7 +147,8 @@ inline constexpr std::string_view kCfgTripleKeys[] = { "arch", "env", "family", "os", }; inline constexpr std::string_view kCfgLayerKeys[] = { - "c++-abi", "c-abi", "compiler", "compiler-runtime", "kernel-abi", + "accelerator", "c++-abi", "c-abi", "compiler", "compiler-runtime", + "kernel-abi", }; inline constexpr std::string_view kCfgBarewords[] = { "linux", "macos", "unix", "windows", @@ -189,7 +211,7 @@ struct Parser { // skipped — which is correct, because the second pass owns it and would // otherwise append the same inputs twice through `append()`. if (is_cfg_layer_key(k)) - return c.layersKnown && c.layer_value(k) == v; + return c.layersKnown && c.layer_matches(k, v); return false; } bool expr() { diff --git a/src/cli.cppm b/src/cli.cppm index 9dfcca6a9..bd1565dd1 100644 --- a/src/cli.cppm +++ b/src/cli.cppm @@ -367,6 +367,10 @@ int run(int argc, char** argv) { .help("Deprecated alias for --cache=off (also clears the build dir)")) .option(cl::Option("target").takes_value().help( "Build for (e.g. x86_64-linux-musl); looks up [target.] in mcpp.toml")) + .option(cl::Option("accel").takes_value().value_name("SPEC") + .help("Device backends and architectures, e.g. 'cuda12.8+{sm_89}'; overrides [build] accel")) + .option(cl::Option("no-accel") + .help("Target no accelerator, ignoring [build] accel")) .option(cl::Option("static").help( "Force static linking (-static). On Linux, prefer pairing with --target -linux-musl")) .option(cl::Option("package").short_name('p').takes_value().value_name("NAME") diff --git a/src/cli/cmd_build.cppm b/src/cli/cmd_build.cppm index a2f55ed9c..18a476cbb 100644 --- a/src/cli/cmd_build.cppm +++ b/src/cli/cmd_build.cppm @@ -90,6 +90,13 @@ export int cmd_build(const mcpplibs::cmdline::ParsedArgs& parsed) { mcpp::build::BuildOverrides ov; if (auto t = parsed.value("target")) ov.target_triple = *t; + // --accel / --no-accel stand to `[build] accel` exactly as --target stands + // to [toolchain]: the manifest declares, the command line overrides for one + // build. --no-accel is not the absence of --accel; it is an explicit + // request for none, which is what a user needs in order to take a CPU-only + // variant of a package that also publishes device builds. + if (parsed.is_flag_set("no-accel")) ov.accel = "(none)"; + else if (auto a = parsed.value("accel")) ov.accel = *a; if (auto p = parsed.value("package")) ov.package_filter = *p; // --cache global|local|off. --no-cache is the deprecated alias for off; the // old flag only ever cleared target/, which says nothing about a cache, so diff --git a/src/doctor.cppm b/src/doctor.cppm index c34ec42dc..da5daa7f6 100644 --- a/src/doctor.cppm +++ b/src/doctor.cppm @@ -34,6 +34,7 @@ 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; @@ -577,6 +578,81 @@ 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())); + } + } + } + } + std::println(""); if (errors) std::println("Doctor result: {} errors, {} warnings", errors, warns); else if (warns) std::println("Doctor result: {} warnings", warns); @@ -1023,6 +1099,17 @@ export int why_report(const std::string& topic) { } std::println(" declared platforms: {} (CI matrix hint)", ps); } + // The declaration, beside the platform one it mirrors. What a build + // actually targets is a different question and a different field + // (`[build] accel` / --accel), reported by the build itself. + if (!ctx->manifest.package.accelerators.empty()) { + std::string as; + for (auto& a : ctx->manifest.package.accelerators) { + if (!as.empty()) as += ", "; + as += a; + } + std::println(" declared accelerators: {} (CI matrix hint)", as); + } } if (all) (void)print_stored_runtime_resolution(); // ⭐ WHERE A NAMED RUNNER BECOMES DISCOVERABLE. diff --git a/src/pack/abi_tag.cppm b/src/pack/abi_tag.cppm index 64322d0eb..447ce5162 100644 --- a/src/pack/abi_tag.cppm +++ b/src/pack/abi_tag.cppm @@ -52,11 +52,34 @@ export namespace mcpp::pack { // A parsed compatibility tag. An EMPTY dimension means "not constrained", // never "unknown": the producer decides what to name, and a consumer must not // invent a constraint the artifact did not declare. +// WHAT DEVICE CODE AN ARTIFACT CARRIES, for one backend. +// +// A PARALLEL FIELD RATHER THAN A TAG SEGMENT, on purpose. The tag is `-` +// joined and parsed from the end, and an architecture list is a SET: joining +// it in would put separators inside a segment of a dash-delimited string whose +// triple already contains a variable number of dashes. A file name is the +// wrong place to carry a set. `tag_check` still compares it, so there is one +// comparator and two storage locations rather than a second comparator. +struct AccelSet { + std::string backend; // "cuda" | "rocm" | … — never empty + std::string version; // toolkit version; empty = unconstrained + std::vector archs; // real device code actually emitted + // The virtual architecture whose portable form is embedded, if any. NVIDIA + // embeds PTX so later hardware can JIT; AMD has no equivalent and obtains + // the same reach through family targets on the `archs` side instead. An + // empty floor must therefore widen nothing. + std::string ptxFloor; +}; + struct AbiTag { std::string triple; // canonical, e.g. "x86_64-linux-gnu" (never empty) std::string compiler; // "gcc16" — empty on a C-surface tag std::string stdlib; // "libstdcxx16" — empty on a C-surface tag std::string standard; // "c++23" — empty on a C-surface tag + // Empty means the artifact carries no device code, which is why a CPU-only + // library is usable by every build: the same don't-care rule as the four + // above, reaching one dimension further. + std::vector accel; bool c_surface() const { return compiler.empty() && stdlib.empty() && standard.empty(); } @@ -120,9 +143,38 @@ struct TagMismatch { // is compared as a floor, not for equality. std::vector tag_check(const AbiTag& published, const AbiTag& current); +// The wire and diagnostic form of the device dimension: +// +// cuda12.8+{sm_80,sm_90f} ptx>=90, rocm6.4+{gfx942} +// +// ONE form for both, so what a descriptor stores is what a refusal prints and +// a reader never has to hold two spellings of the same fact in their head. +// The backend is the leading run of letters and the version is what follows, +// which parses because every backend name is alphabetic and every version +// starts with a digit. +std::string accel_str(std::span sets); + +// The inverse. Unparseable input yields an empty vector, which the comparison +// reads as "carries no device code" — the same answer as a descriptor that +// never mentioned the dimension, and the safe one. +std::vector parse_accel(std::string_view s); + // The `c++NN` segment as its numeric level, or 0 when unparseable. int standard_level(std::string_view standardSegment); +// The numeric level of a device architecture, or 0 when it is not of the +// numbered `sm_`/`compute_` shape. AMD's `gfx942` answers 0 on purpose: it +// carries no ordering that means anything here, so it compares by equality. +int accel_arch_level(std::string_view arch); + +// Does the device code `published` names cover `wantedArch`? +bool accel_arch_covers(std::string_view published, std::string_view wantedArch); + +// Does `published` satisfy every backend and architecture `wanted` asks for? +// An empty `published` is unconstrained; an empty `wanted` is satisfied. +bool accel_accepts(std::span published, + std::span wanted); + } // namespace mcpp::pack namespace mcpp::pack { @@ -167,6 +219,155 @@ AbiTag cxx_surface_tag(const mcpp::toolchain::Toolchain& tc, return t; } +std::string accel_str(std::span sets) { + if (sets.empty()) return "(none)"; + std::string out; + for (auto const& a : sets) { + if (!out.empty()) out += ", "; + out += a.backend; + out += a.version; + out += "+{"; + for (std::size_t i = 0; i < a.archs.size(); ++i) { + if (i) out += ','; + out += a.archs[i]; + } + out += '}'; + if (!a.ptxFloor.empty()) { out += " ptx>="; out += a.ptxFloor; } + } + return out; +} + +namespace { + +std::string_view trim_sv(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; +} + +} // namespace + +std::vector parse_accel(std::string_view s) { + std::vector out; + for (std::size_t i = 0; i <= s.size(); ) { + auto comma = s.find(',', i); + // A comma inside `{...}` separates architectures, not backends. + auto open = s.find('{', i); + auto close = s.find('}', i); + if (open != std::string_view::npos && close != std::string_view::npos + && comma != std::string_view::npos && comma > open && comma < close) + comma = s.find(',', close); + auto chunk = trim_sv(comma == std::string_view::npos + ? s.substr(i) + : s.substr(i, comma - i)); + i = comma == std::string_view::npos ? s.size() + 1 : comma + 1; + if (chunk.empty() || chunk == "(none)") continue; + + AccelSet a; + auto plus = chunk.find('+'); + auto head = trim_sv(plus == std::string_view::npos ? chunk : chunk.substr(0, plus)); + std::size_t n = 0; + while (n < head.size() && std::isalpha(static_cast(head[n]))) ++n; + if (n == 0) continue; // no backend: not our form + a.backend = std::string(head.substr(0, n)); + a.version = std::string(head.substr(n)); + + if (plus != std::string_view::npos) { + auto tail = chunk.substr(plus + 1); + if (auto o = tail.find('{'); o != std::string_view::npos) { + auto c = tail.find('}', o); + auto archs = tail.substr(o + 1, + c == std::string_view::npos ? std::string_view::npos : c - o - 1); + for (std::size_t j = 0; j <= archs.size(); ) { + auto k = archs.find(',', j); + auto one = trim_sv(k == std::string_view::npos + ? archs.substr(j) : archs.substr(j, k - j)); + if (!one.empty()) a.archs.emplace_back(one); + if (k == std::string_view::npos) break; + j = k + 1; + } + } + if (auto pf = tail.find("ptx>="); pf != std::string_view::npos) + a.ptxFloor = std::string(trim_sv(tail.substr(pf + 5))); + } + out.push_back(std::move(a)); + } + return out; +} + +int accel_arch_level(std::string_view arch) { + auto digits = arch; + if (digits.starts_with("sm_")) digits.remove_prefix(3); + else if (digits.starts_with("compute_")) digits.remove_prefix(8); + int level = 0, n = 0; + for (char c : digits) { + if (!std::isdigit(static_cast(c))) break; + level = level * 10 + (c - '0'); + ++n; + } + return n == 0 ? 0 : level; +} + +namespace { + +// The trailing letter of a qualified target, or '\0'. +// +// sm_90 baseline that compute capability +// sm_90f family same major, equal-or-higher minor +// sm_90a architecture that compute capability and no other +char accel_arch_suffix(std::string_view arch) { + if (arch.empty()) return '\0'; + char last = arch.back(); + return (last == 'f' || last == 'a') ? last : '\0'; +} + +} // namespace + +bool accel_arch_covers(std::string_view published, std::string_view wantedArch) { + if (published == wantedArch) return true; + int have = accel_arch_level(published); + int want = accel_arch_level(wantedArch); + // Not the numbered shape on either side: equality was the only question + // available, and it has been answered. + if (have == 0 || want == 0) return false; + if (accel_arch_suffix(published) == 'f') + return have / 10 == want / 10 && want % 10 >= have % 10; + // Baseline and architecture-specific targets cover exactly their own + // compute capability. Forward reach for those comes from embedded portable + // code, which is ptxFloor's question rather than this one. + return have == want; +} + +bool accel_accepts(std::span published, + std::span wanted) +{ + if (published.empty()) return true; // carries no device code + for (auto const& want : wanted) { + const AccelSet* have = nullptr; + for (auto const& p : published) + if (p.backend == want.backend) { have = &p; break; } + if (!have) return false; + // Minor version compatibility is real inside one major release family + // and absent across families, so the major is compared and the minor + // is not. + if (!have->version.empty() && !want.version.empty() + && major_of(have->version) != major_of(want.version)) + return false; + int floor = have->ptxFloor.empty() ? 0 : accel_arch_level(have->ptxFloor); + for (auto const& a : want.archs) { + bool ok = false; + for (auto const& p : have->archs) + if (accel_arch_covers(p, a)) { ok = true; break; } + if (!ok && floor != 0) { + int lvl = accel_arch_level(a); + ok = lvl != 0 && lvl >= floor; + } + if (!ok) return false; + } + } + return true; +} + int standard_level(std::string_view seg) { if (!seg.starts_with("c++")) return 0; int level = 0; @@ -225,6 +426,16 @@ std::vector tag_check(const AbiTag& published, const AbiTag& curren if (need != 0 && got != 0 && got < need) out.push_back({ "standard", published.standard, current.standard }); } + + // The device dimension. Reported as ONE mismatch rather than one per + // backend: a consumer cannot act on "cuda disagrees and rocm disagrees" + // any differently than on "this artifact's device code does not cover the + // target", and the closest-refusal selection in prebuilt.cppm ranks + // candidates by how many dimensions disagree. + if (!accel_accepts(published.accel, current.accel)) + out.push_back({ "accel", + accel_str(published.accel), + accel_str(current.accel) }); return out; } diff --git a/src/pack/prebuilt.cppm b/src/pack/prebuilt.cppm index 0e6e5b899..50a0b6a14 100644 --- a/src/pack/prebuilt.cppm +++ b/src/pack/prebuilt.cppm @@ -104,9 +104,16 @@ check_prebuilt(const mcpp::manifest::Manifest& m, const PrebuiltCheck& in) accepted = true; continue; } - publishedTags.push_back(a.abi); + // The device dimension is not part of the tag string, so a listing + // that printed only the tag would show two indistinguishable lines for + // a package publishing a CPU and a device build of the same triple. + publishedTags.push_back(a.accel.empty() + ? a.abi + : std::format("{} accel={}", a.abi, a.accel)); auto published = parse_abi_tag(a.abi); if (!published) { accepted = true; continue; } // unreadable → lenient + // The device dimension travels beside the tag rather than inside it. + published->accel = parse_accel(a.accel); auto bad = tag_check(*published, in.current); if (bad.empty()) { accepted = true; break; } // Keep the CLOSEST refusal to show: the one that disagrees least is @@ -121,16 +128,35 @@ check_prebuilt(const mcpp::manifest::Manifest& m, const PrebuiltCheck& in) std::string tags; for (auto const& t : publishedTags) tags += std::format("\n {}", t); std::string why; - for (auto const& b : bestRefusal) + bool accelOnly = !bestRefusal.empty(); + for (auto const& b : bestRefusal) { why += std::format("\n {:<9} needs {}, this build has {}", b.dimension, b.need, b.got); + if (b.dimension != "accel") accelOnly = false; + } + // The remedy differs by dimension, and the generic one is actively + // misleading for the device axis: pinning [toolchain] cannot change + // which GPU architecture a build targets. + const std::string fix = accelOnly + ? std::string( + " fix: build for an architecture the package carries (--accel), or take\n" + " a variant that carries no device code (--no-accel), or ask the\n" + " publisher for one covering yours.") + : std::string( + " fix: ask the publisher for a build matching your toolchain, or pin\n" + " [toolchain] to the one the package was built with."); + // The device dimension travels beside the tag, so the summary line has + // to name it too or it contradicts the per-dimension list below. + const std::string current = in.current.accel.empty() + ? in.current.str() + : std::format("{} accel={}", in.current.str(), + accel_str(in.current.accel)); return std::unexpected(std::format( "{}: no prebuilt artifact matches this toolchain.\n" " your toolchain : {}\n" " published tags :{}\n" " closest is {}, and it differs on:{}\n" - " fix: ask the publisher for a build matching your toolchain, or pin\n" - " [toolchain] to the one the package was built with.", - in.packageLabel, in.current.str(), tags, bestRefusalTag, why)); + "{}", + in.packageLabel, current, tags, bestRefusalTag, why, fix)); } // ── 3. the interface is the one the binaries were built from ────── diff --git a/src/toolchain/devicehost.cppm b/src/toolchain/devicehost.cppm new file mode 100644 index 000000000..3a4626084 --- /dev/null +++ b/src/toolchain/devicehost.cppm @@ -0,0 +1,100 @@ +// 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); + +// 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; +} + +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/tests/e2e/600_accel_variant_selection.sh b/tests/e2e/600_accel_variant_selection.sh new file mode 100755 index 000000000..bd4665919 --- /dev/null +++ b/tests/e2e/600_accel_variant_selection.sh @@ -0,0 +1,114 @@ +#!/usr/bin/env bash +# requires: elf +# The accelerator dimension decides which prebuilt artifact a build may take. +# +# Two packages, because the two halves of the rule need different shapes: +# +# gpuonly publishes ONE artifact, for sm_90f. A build targeting sm_86 has +# nothing it can use, so it must be refused — and the refusal must +# name the dimension, not a digest or a triple. +# gpukit publishes a CPU-only artifact FIRST and a device one after it. +# The same sm_86 build must succeed by taking the first, which is +# the ordering that also keeps an mcpp predating this field working. +set -e + +TMP=$(mktemp -d) +trap "rm -rf $TMP" EXIT +cd "$TMP" + +TRIPLE="x86_64-linux-gnu" + +make_pkg() { # $1 name, $2... artifact tables + local name="$1"; shift + mkdir -p "$name/lib" "$name/include" + : > "$name/lib/lib${name}.a" + { + echo '[package]' + echo "name = \"$name\"" + echo 'version = "0.1.0"' + echo '[language]' + echo 'standard = "c++23"' + printf '%s\n' "$@" + } > "$name/mcpp.toml" +} + +# One device-only artifact. +make_pkg gpuonly \ +'[[runtime.artifacts]]' \ +'role = "static-library"' \ +'path = "lib/libgpuonly.a"' \ +'provenance = "mcpp-pack/1"' \ +"abi = \"$TRIPLE\"" \ +'accel = "cuda12.8+{sm_90f}"' + +# CPU-only first, device second. +make_pkg gpukit \ +'[[runtime.artifacts]]' \ +'role = "static-library"' \ +'path = "lib/libgpukit.a"' \ +'provenance = "mcpp-pack/1"' \ +"abi = \"$TRIPLE\"" \ +'' \ +'[[runtime.artifacts]]' \ +'role = "static-library"' \ +'path = "lib/libgpukit.a"' \ +'provenance = "mcpp-pack/1"' \ +"abi = \"$TRIPLE\"" \ +'accel = "cuda12.8+{sm_90f}"' + +consumer() { # $1 dep name + rm -rf app; "$MCPP" new app > /dev/null; cd app + cat > mcpp.toml < src/main.cpp <<'EOF' +int main() { return 0; } +EOF + rm -f src/*.cppm + cd .. +} + +# ── 1. a device-only package refuses a build targeting another architecture ── +consumer gpuonly +cd app +if "$MCPP" build --accel 'cuda12.8+{sm_86}' > out.log 2>&1; then + cat out.log; echo "FAIL: sm_86 accepted an sm_90f-only artifact"; exit 1 +fi +grep -q 'accel' out.log || { cat out.log; echo "FAIL: refusal does not name the accel dimension"; exit 1; } +grep -q 'sm_90f' out.log || { cat out.log; echo "FAIL: refusal does not say what the artifact has"; exit 1; } +grep -q 'sm_86' out.log || { cat out.log; echo "FAIL: refusal does not say what this build asked for"; exit 1; } +# The remedy has to be the one that can actually work. Pinning [toolchain] +# cannot change which GPU architecture a build targets, so the generic advice +# is worse than none here. +grep -q -- '--no-accel' out.log \ + || { cat out.log; echo "FAIL: refusal offers no remedy for the device axis"; exit 1; } +grep -q 'accel=cuda12.8+{sm_90f}' out.log \ + || { cat out.log; echo "FAIL: the published listing hides the device dimension"; exit 1; } +cd .. + +# ── 2. the same package accepts the architecture it was built for ──────────── +cd app +"$MCPP" build --accel 'cuda12.8+{sm_90}' > ok.log 2>&1 \ + || { cat ok.log; echo "FAIL: sm_90 refused by an sm_90f artifact"; exit 1; } +cd .. + +# ── 3. a build asking for no accelerator is satisfied vacuously ────────────── +cd app +"$MCPP" build --no-accel > none.log 2>&1 \ + || { cat none.log; echo "FAIL: --no-accel refused"; exit 1; } +cd .. + +# ── 4. CPU-first ordering lets an unmatched device build fall back ─────────── +consumer gpukit +cd app +"$MCPP" build --accel 'cuda12.8+{sm_86}' > fallback.log 2>&1 \ + || { cat fallback.log; echo "FAIL: CPU-only artifact listed first did not accept sm_86"; exit 1; } +cd .. + +echo "PASS: accel variant selection" diff --git a/tests/unit/test_abi_tag_accel.cpp b/tests/unit/test_abi_tag_accel.cpp new file mode 100644 index 000000000..a06486702 --- /dev/null +++ b/tests/unit/test_abi_tag_accel.cpp @@ -0,0 +1,161 @@ +#include + +import std; +import mcpp.pack.abi_tag; + +using mcpp::pack::AbiTag; +using mcpp::pack::AccelSet; +using mcpp::pack::tag_check; + +namespace { + +AbiTag cxx() { + AbiTag t; + t.triple = "x86_64-linux-gnu"; t.compiler = "gcc16"; + t.stdlib = "libstdcxx16"; t.standard = "c++23"; + return t; +} + +AbiTag with(std::vector a) { auto t = cxx(); t.accel = std::move(a); return t; } + +AccelSet cuda(std::vector archs, std::string ptx = {}) { + return AccelSet{ .backend = "cuda", .version = "12.8", + .archs = std::move(archs), .ptxFloor = std::move(ptx) }; +} + +bool accepts(const AbiTag& published, const AbiTag& current) { + return tag_check(published, current).empty(); +} + +} // namespace + +// ─── The published side constrains; an absent dimension does not ─────────── + +TEST(AbiTagAccel, PublishedWithoutAccelAcceptsAnyRequest) { + // A CPU-only library states nothing about accelerators, so it is usable + // everywhere. This is the existing "a shorter tag IS the statement" rule + // reaching one dimension further. + EXPECT_TRUE(accepts(cxx(), with({cuda({"sm_86"})}))); +} + +TEST(AbiTagAccel, ARequestWithNoAcceleratorIsSatisfiedVacuously) { + EXPECT_TRUE(accepts(with({cuda({"sm_90"})}), cxx())); +} + +// ─── Membership ──────────────────────────────────────────────────────────── + +TEST(AbiTagAccel, ExactArchIsAccepted) { + EXPECT_TRUE(accepts(with({cuda({"sm_80", "sm_90"})}), with({cuda({"sm_90"})}))); +} + +TEST(AbiTagAccel, AnArchOutsideTheSetIsRefused) { + auto bad = tag_check(with({cuda({"sm_80", "sm_90"})}), with({cuda({"sm_86"})})); + ASSERT_EQ(bad.size(), 1u); + EXPECT_EQ(bad[0].dimension, "accel"); + EXPECT_NE(bad[0].need.find("sm_80"), std::string::npos); + EXPECT_NE(bad[0].got.find("sm_86"), std::string::npos); +} + +TEST(AbiTagAccel, ADifferentBackendIsRefused) { + AccelSet rocm{ .backend = "rocm", .version = "6.4", .archs = {"gfx942"} }; + auto bad = tag_check(with({rocm}), with({cuda({"sm_90"})})); + ASSERT_EQ(bad.size(), 1u); + EXPECT_NE(bad[0].need.find("rocm"), std::string::npos); +} + +// ─── Family targets widen the published side ─────────────────────────────── + +TEST(AbiTagAccel, AFamilyTargetCoversItsFamily) { + // sm_90f is compatible with the same major and an equal-or-higher minor. + // This is why publishing a family target rather than one artifact per chip + // is what keeps the variant matrix finite. + EXPECT_TRUE (accepts(with({cuda({"sm_90f"})}), with({cuda({"sm_90"})}))); + EXPECT_TRUE (accepts(with({cuda({"sm_90f"})}), with({cuda({"sm_93"})}))); + EXPECT_FALSE(accepts(with({cuda({"sm_90f"})}), with({cuda({"sm_89"})}))); + EXPECT_FALSE(accepts(with({cuda({"sm_90f"})}), with({cuda({"sm_100"})}))); +} + +TEST(AbiTagAccel, AnArchitectureSpecificTargetCoversOnlyItself) { + EXPECT_TRUE (accepts(with({cuda({"sm_90a"})}), with({cuda({"sm_90"})}))); + EXPECT_FALSE(accepts(with({cuda({"sm_90a"})}), with({cuda({"sm_93"})}))); +} + +// ─── The PTX floor is the second asymmetric dimension ────────────────────── + +TEST(AbiTagAccel, EmbeddedPtxAcceptsAnythingAtOrAboveTheFloor) { + // NVIDIA's own guidance for a distributed binary is cubins for the known + // targets plus PTX for the newest, so that later hardware JITs. Recording + // the floor is what lets tag_check say yes to a device the cubin set does + // not name, instead of refusing an artifact that genuinely runs. + EXPECT_TRUE (accepts(with({cuda({"sm_80"}, "80")}), with({cuda({"sm_86"})}))); + EXPECT_TRUE (accepts(with({cuda({"sm_80"}, "80")}), with({cuda({"sm_90"})}))); + EXPECT_FALSE(accepts(with({cuda({"sm_86"}, "86")}), with({cuda({"sm_80"})}))); +} + +TEST(AbiTagAccel, WithoutPtxThereIsNoForwardCompatibility) { + // AMD has no PTX equivalent, so an empty floor must not silently widen the + // set. Family coverage on the published side is how ROCm gets the same + // effect, and that is a different mechanism. + EXPECT_FALSE(accepts(with({cuda({"sm_80"})}), with({cuda({"sm_86"})}))); +} + +// ─── Toolkit version ─────────────────────────────────────────────────────── + +TEST(AbiTagAccel, TheToolkitMajorMustAgree) { + AccelSet built{ .backend = "cuda", .version = "12.8", .archs = {"sm_90"} }; + AccelSet want12{ .backend = "cuda", .version = "12.0", .archs = {"sm_90"} }; + AccelSet want13{ .backend = "cuda", .version = "13.0", .archs = {"sm_90"} }; + // Minor version compatibility is real within a major release family. + EXPECT_TRUE (accepts(with({built}), with({want12}))); + EXPECT_FALSE(accepts(with({built}), with({want13}))); +} + +// ─── Several backends in one artifact ────────────────────────────────────── + +TEST(AbiTagAccel, OneArtifactCanCarrySeveralBackends) { + AccelSet rocm{ .backend = "rocm", .version = "6.4", .archs = {"gfx942"} }; + auto fat = with({cuda({"sm_90f"}), rocm}); + EXPECT_TRUE(accepts(fat, with({cuda({"sm_90"})}))); + AccelSet wantRocm{ .backend = "rocm", .version = "6.4", .archs = {"gfx942"} }; + EXPECT_TRUE(accepts(fat, with({wantRocm}))); +} + +// ─── The other four dimensions still decide ──────────────────────────────── + +TEST(AbiTagAccel, AccelDoesNotMaskTheCxxDimensions) { + auto published = with({cuda({"sm_90"})}); + auto current = with({cuda({"sm_90"})}); + current.compiler = "gcc15"; + auto bad = tag_check(published, current); + ASSERT_EQ(bad.size(), 1u); + EXPECT_EQ(bad[0].dimension, "compiler"); +} + +// ─── The wire form round-trips ───────────────────────────────────────────── + +TEST(AbiTagAccel, TheWireFormRoundTrips) { + std::vector sets{ + cuda({"sm_80", "sm_90f"}, "90"), + mcpp::pack::AccelSet{ .backend = "rocm", .version = "6.4", + .archs = {"gfx942", "gfx10-3-generic"} }, + }; + auto text = mcpp::pack::accel_str(sets); + auto back = mcpp::pack::parse_accel(text); + ASSERT_EQ(back.size(), 2u); + EXPECT_EQ(back[0].backend, "cuda"); + EXPECT_EQ(back[0].version, "12.8"); + EXPECT_EQ(back[0].archs, (std::vector{"sm_80", "sm_90f"})); + EXPECT_EQ(back[0].ptxFloor, "90"); + EXPECT_EQ(back[1].backend, "rocm"); + EXPECT_EQ(back[1].archs, (std::vector{"gfx942", "gfx10-3-generic"})); + EXPECT_TRUE(back[1].ptxFloor.empty()); + EXPECT_EQ(mcpp::pack::accel_str(back), text); +} + +TEST(AbiTagAccel, UnparseableTextMeansNoDeviceCode) { + // The safe answer, and the same one a descriptor that never mentioned the + // dimension gives: an artifact that states nothing constrains nothing. + EXPECT_TRUE(mcpp::pack::parse_accel("").empty()); + EXPECT_TRUE(mcpp::pack::parse_accel("(none)").empty()); + EXPECT_TRUE(mcpp::pack::parse_accel("12.8+{sm_90}").empty()); // no backend +} diff --git a/tests/unit/test_cfg_accelerator.cpp b/tests/unit/test_cfg_accelerator.cpp new file mode 100644 index 000000000..f6f3a1afc --- /dev/null +++ b/tests/unit/test_cfg_accelerator.cpp @@ -0,0 +1,86 @@ +#include + +import std; +import mcpp.build.prepare_inputs; + +namespace cfgpred = mcpp::build::cfgpred; + +namespace { + +// A context with the triple coordinates filled and the target side resolved, +// which is the state the second merge pass runs in. +cfgpred::Ctx resolved(std::vector accelerators) { + auto c = cfgpred::context_for("x86_64-linux-gnu"); + c.layersKnown = true; + c.accelerators = std::move(accelerators); + return c; +} + +bool m(std::string_view predicate, const cfgpred::Ctx& c) { + return cfgpred::matches(std::string(predicate), c); +} + +} // namespace + +// ─── Membership, not equality ────────────────────────────────────────────── +// +// `accelerator` is the first MULTI-VALUED layer: one build can enable several +// backends at once, which is what an inference framework shipping CUDA and +// ROCm in one artifact requires. The comparison is therefore membership. +// +// The alternative considered and rejected was to let `any(...)` mean +// membership while a bare key meant set equality. That would make a +// combinator change the meaning of its operand, and `all(a = "cuda", +// a = "rocm")` would then be unsatisfiable rather than "both are enabled". +// Membership everywhere keeps `any`/`all`/`not` pure boolean combinators. + +TEST(CfgAccelerator, BareKeyIsMembership) { + auto cuda = resolved({"cuda"}); + EXPECT_TRUE (m(R"(cfg(accelerator = "cuda"))", cuda)); + EXPECT_FALSE(m(R"(cfg(accelerator = "rocm"))", cuda)); +} + +TEST(CfgAccelerator, MultipleBackendsAreAllMembers) { + auto both = resolved({"cuda", "rocm"}); + EXPECT_TRUE(m(R"(cfg(accelerator = "cuda"))", both)); + EXPECT_TRUE(m(R"(cfg(accelerator = "rocm"))", both)); +} + +TEST(CfgAccelerator, CombinatorsStayPure) { + auto both = resolved({"cuda", "rocm"}); + auto cuda = resolved({"cuda"}); + EXPECT_TRUE (m(R"(cfg(any(accelerator = "cuda", accelerator = "rocm")))", cuda)); + EXPECT_TRUE (m(R"(cfg(all(accelerator = "cuda", accelerator = "rocm")))", both)); + EXPECT_FALSE(m(R"(cfg(all(accelerator = "cuda", accelerator = "rocm")))", cuda)); + EXPECT_TRUE (m(R"(cfg(not(accelerator = "rocm")))", cuda)); +} + +TEST(CfgAccelerator, ComposesWithTripleKeys) { + auto cuda = resolved({"cuda"}); + EXPECT_TRUE (m(R"(cfg(all(linux, accelerator = "cuda")))", cuda)); + EXPECT_FALSE(m(R"(cfg(all(windows, accelerator = "cuda")))", cuda)); +} + +TEST(CfgAccelerator, UnresolvedTargetSideDoesNotMatch) { + // The first pass runs before dependency resolution and cannot answer a + // layer key. Returning false there is correct: the second pass owns it and + // would otherwise contribute the same inputs twice. + auto early = cfgpred::context_for("x86_64-linux-gnu"); + ASSERT_FALSE(early.layersKnown); + EXPECT_FALSE(m(R"(cfg(accelerator = "cuda"))", early)); +} + +TEST(CfgAccelerator, IsALayerKeyNotATripleKey) { + EXPECT_TRUE (cfgpred::uses_layer(R"(cfg(accelerator = "cuda"))")); + EXPECT_FALSE(cfgpred::uses_layer(R"(cfg(arch = "x86_64"))")); +} + +TEST(CfgAccelerator, MisspellingIsReportedNotSilentlyFalse) { + // A predicate naming a token outside the vocabulary used to evaluate to + // false in silence, which is indistinguishable from one that correctly did + // not apply. + auto unknown = cfgpred::unknown_tokens(R"(cfg(acclerator = "cuda"))"); + ASSERT_EQ(unknown.size(), 1u); + EXPECT_EQ(unknown[0], "acclerator"); + EXPECT_TRUE(cfgpred::unknown_tokens(R"(cfg(accelerator = "cuda"))").empty()); +} diff --git a/tests/unit/test_devicehost.cpp b/tests/unit/test_devicehost.cpp new file mode 100644 index 000000000..9d7b32bee --- /dev/null +++ b/tests/unit/test_devicehost.cpp @@ -0,0 +1,58 @@ +#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)); +} diff --git a/tests/unit/test_source_kind.cpp b/tests/unit/test_source_kind.cpp index 2e5e1f721..b49c2b6d2 100644 --- a/tests/unit/test_source_kind.cpp +++ b/tests/unit/test_source_kind.cpp @@ -225,3 +225,67 @@ TEST(SourceKind, ModuleInterfaceLangFlagIsPerCompilerAndNotInterchangeable) { mcpp::toolchain::CompilerId::MSVC}) EXPECT_FALSE(traits_for(id).moduleInterfaceLangFlag.empty()); } + +// ─── T-6. Device translation units ───────────────────────────────────────── +// +// A device TU is compiled by a vendor compiler (nvcc, hipcc) that mcpp does +// not drive directly, and no such compiler accepts C++20 modules. The kind +// therefore states the GRAPH ROLE — "not scanned, no BMI" — and says nothing +// about the language, which is what lets one kind cover CUDA C++, HIP, and +// dialects that are not C++ at all. + +TEST(SourceKind, ClassifiesDeviceTranslationUnits) { + auto t = mcpp::builtin_extension_table(); + EXPECT_EQ(mcpp::classify("src/k.cu", t), SourceKind::Device); + EXPECT_EQ(mcpp::classify("src/k.hip", t), SourceKind::Device); +} + +TEST(SourceKind, DeviceHeadersAffectGraphShape) { + // `.cuh` reaches a device TU through the preprocessor, so editing one can + // change what the graph should be. Leaving it in `Other` is why a project + // would see "edited the kernel header, nothing rebuilt". + auto t = mcpp::builtin_extension_table(); + EXPECT_EQ(mcpp::classify("src/k.cuh", t), SourceKind::Header); + EXPECT_EQ(mcpp::classify("src/k.hiph", t), SourceKind::Header); + EXPECT_TRUE(mcpp::affects_graph_shape(SourceKind::Header)); +} + +TEST(SourceKind, DevicePredicates) { + // Scan-exempt for the same reason assembly is: a device TU carries no + // `import`, so there is nothing for P1689 to answer. + EXPECT_TRUE (mcpp::is_scan_exempt(SourceKind::Device)); + EXPECT_FALSE(mcpp::produces_bmi(SourceKind::Device)); + EXPECT_FALSE(mcpp::is_cxx_like(SourceKind::Device)); + EXPECT_FALSE(mcpp::links_unconditionally(SourceKind::Device)); + // Absent for the same reason assembly is absent: the content change is + // tracked by ninja, and a NEW file is `glob_inputs_stale`'s question. + EXPECT_FALSE(mcpp::affects_graph_shape(SourceKind::Device)); + EXPECT_EQ(mcpp::to_string(SourceKind::Device), "device"); +} + +TEST(SourceKind, DeviceExtensionsAreNotInTheDefaultGlob) { + // Same compatibility argument as the built-in module-extension table: a + // published package that vendors a `.cu` it builds elsewhere must not + // start compiling it on the next mcpp upgrade. Device sources are opted + // into by naming them in a `kind = "device"` target. + auto globs = mcpp::default_source_globs(mcpp::builtin_extension_table()); + for (auto const& g : globs) { + EXPECT_EQ(g.find(".cu"), std::string::npos) << g; + EXPECT_EQ(g.find(".hip"), std::string::npos) << g; + } + EXPECT_EQ(mcpp::default_source_globs_note(mcpp::builtin_extension_table()) + .find("cu"), std::string::npos); +} + +TEST(SourceKind, DeviceObjectsUseTheCollisionProofName) { + // `k.cu` and `k.cpp` in one directory are common in a mixed project; the + // stem-named form would give both `k.o`. + EXPECT_EQ(mcpp::object_filename_for("src/k.cu", ".o"), "k.cu.o"); + EXPECT_EQ(mcpp::object_filename_for("src/k.hip", ".o"), "k.hip.o"); +} + +TEST(SourceKind, DeviceExtensionsCannotBeClaimedAsModuleInterfaces) { + auto err = mcpp::validate_module_extensions(std::vector{".cu"}); + ASSERT_TRUE(err.has_value()); + EXPECT_NE(err->find(".cu"), std::string::npos); +}