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

* feat(doctor): read the device toolkit from a payload before the host, and state the driver relation without probing for it

## The report was answering about the wrong toolkit

`mcpp self doctor` located `crt/host_config.h` through `CUDA_PATH`, `CUDA_HOME`,
`/usr/local/cuda` and `/usr/include` — host locations only. A toolkit installed
through xlings is the one a build will use, and it is usually the newer one:
measured here, a 12.9 payload states `gcc <= 14` and a 13.3 payload `gcc <= 15`,
where this machine's distribution CUDA 12.0 states `gcc <= 12`. The report said
`gcc 13 exceeds the bound of 12` about a toolkit the build was not using.

Payload stores are now searched first. Both of them: mcpp keeps its own under
`<mcpp home>/registry/data/xpkgs`, and `xlings install` writes to
`<xlings home>/data/xpkgs` — 191 packages in one and 211 in the other on this
machine, with the CUDA components only in the second. The host locations remain,
last, because a machine with a distribution toolkit and no payload is real.

Reading changes accordingly:

    before   warning: cuda will refuse this host compiler: gcc 13 exceeds
                      the bound of 12 stated in /usr/include/crt/host_config.h
    after    ok       cuda accepts this host compiler (gcc 13 <= 15)

## The driver relation, and why the acquisition is not here

A device runtime must not be newer than the driver it runs against. Measured on
a driver serving CUDA 12.4: the 13.3 payload compiles and links cleanly and then
fails at the first allocation with "CUDA driver version is insufficient for CUDA
runtime version", while the 12.9 payload prints the right answer.

`mcpp::toolchain::driver_accepts_toolkit` states when one version may meet
another, including that minor-version compatibility makes 12.9 fine against a
driver serving 12.4 — the case a naive "toolkit <= driver" check would have
refused. Five unit tests, including that either side unknown makes no claim.

The acquisition is NOT here. Asking a machine which driver it has means running
a vendor's tool, and `tests/unit/test_runtime_contract.cpp` refuses exactly that
in `src/`. It caught the first revision of this change, which launched one. The
rule predates this work and it is right: a core that learns to run one vendor's
probe learns to run four. Those numbers will reach the report as declarations —
a toolkit payload stating the driver it needs, and the package that owns the
host driver stating what the host has — which is the rule-package channel.

## Verified

`tests/e2e/602_device_toolkit_payload_first.sh` fabricates a payload store whose
header states `gcc <= 41`, a bound nothing real would state, and asserts the
report reads it. The control is the half that matters: without the payload store
the same command must not report 41, or the assertion would pass against a
doctor that hardcoded it.

100 test binaries pass, including the contract test that rejected the earlier
revision.

* feat(example): the CUDA example takes its toolkit from a payload, not from the host

The example shipped in 2026.9.5.1 wrote `-L/usr/local/cuda/lib64` in its
manifest and probed `/usr/local/cuda/bin/nvcc` and `/usr/include/crt/host_config.h`
in its rule package. It was the project's own demonstration of how to do this,
and what it demonstrated was reaching for the host.

The project now names the toolkit:

    [xlings.workspace]
    "xim:cuda-nvcc"   = "12.9.86"
    "xim:cuda-cudart" = "12.9.79"

and the rule resolves it with `mcpp::xpkg_dir`, building the whole invocation
from what it finds — compiler, include directories, and library search paths
through `mcpp::link_search`. The manifest names libraries and no locations.
Host paths remain in the rule as a last fallback so a machine with only a
distribution toolkit still builds.

⚠️ THE PAYLOAD'S HEADERS HAVE TO BE NAMED. nvcc adds
`<its own directory>/../include` by itself, and on the 12.x line that holds
`crt/` but not `cuda_runtime.h` — which is in the `cuda-cudart` component. The
first revision of this change left it out: nvcc resolved `cuda_runtime.h` from
/usr/include and then read the HOST's `crt/host_config.h` beside it, and the
build failed with the host toolkit's complaint (`clang version must be less than
15`) while using the payload's compiler. Naming the payload include directories
is what makes "uses the payload" true rather than nearly true.

Verified: `mcpp run` prints `12 24 36 48`, and

    $ mcpp build -v | grep -c '/usr/local/cuda\|/usr/bin/nvcc|-I/usr/include'
    0

nvcc and both include directories resolve under
`registry/data/xpkgs/xim-x-cuda-*`. 100 test binaries pass.

* feat(build): version floors, so a machine's limits are known before anything is compiled

Some facts about a machine bound what may be built for it, and the failure when
they are ignored arrives late. Measured on a driver serving CUDA 12.4: a binary
built with the 13.3 toolkit compiles cleanly, links cleanly, and fails at the
first allocation with "CUDA driver version is insufficient for CUDA runtime
version" — a message naming neither the toolkit nor the driver.

Both numbers are knowable beforehand. What mcpp must not do is go and ask a
vendor's tool for them, and `tests/unit/test_runtime_contract.cpp` refuses
exactly that in `src/`. So the numbers arrive as declarations.

A package states what it needs:

    [[runtime.requirements]]
    kind  = "version-floor"
    value = "cuda.driver >= 12.0"

A package that established a fact about this machine — at install time, which
is where probing belongs — states it:

    [runtime]
    provides = ["cuda.driver=12.4"]

They are compared when capabilities are bound, and a short machine is refused
before anything is compiled, reporting `version-floor-unmet`.

## No vendor vocabulary reaches the engine

`src/build/version_floor.cppm` reads a name, a relation and a version and knows
nothing about what any of them mean. `cuda.driver` is data passing through, no
new manifest keys were needed — `kind` was already a free string — and a
backend mcpp has never heard of compares the same way. That is asserted with a
name no backend uses.

## A floor nobody answered is silent

A machine that never declared what it has is not a machine that fails the floor;
it is one nobody asked. Turning "we do not know" into "no" is the failure mode
this exists to avoid, so it is the third case in the e2e rather than a comment.

Without that control the first case would also pass against an engine that
refused every version-floor requirement.

## Verified

`tests/e2e/603_version_floor.sh`, three halves:

    PASS: refused, naming the requirement, both versions and who stated the fact
    PASS: a met floor builds
    PASS: a floor with no stated fact is silent, for a name no backend uses

Eight unit tests over the parsing and the comparison, including that `>=` is
never read as `=` — both spellings live in string lists, and reading a floor as
a fact would turn a requirement into its own satisfaction. 101 test binaries
pass.

* test(e2e): give 317 the window its bound needs, and run 602 where the report exists

Reaching "five consecutive short failures" costs at least 1.25 s (four
restart delays of 250 ms and five 50 ms polls). A two-second pause left
150 ms per spawn, and a loaded macOS runner exceeded it twice in a row on
main while the same code passed twice on this branch. The pause is now
five seconds, so the property is decided by the code.

602 declared no platform, and on Windows the doctor does not emit the
device-toolkit section at all: the payloads are Linux builds and the bound
a Windows toolkit states is an _MSC_VER range the report does not read.
It also ran the doctor in a fresh home, which provisioned a full toolchain
into the temporary directory: 229 s and 1.4 GB per run. It now requires
unix-shell and runs under MCPP_OFFLINE.

* fix(config): offline mode skips the first-use sandbox bootstrap

--offline promises never to touch the network. In a home that had never
been used, load_or_init cloned the package index through `xlings self
init` and installed ninja and patchelf through `xlings install` before any
command ran: measured, 26 s and 126 MB under MCPP_OFFLINE=1.

The three network-bound steps are now skipped under offline mode, the skip
is announced once per process, and the completion markers are still
written from what is on disk, so a bootstrapped home loses nothing. The
commands that need the tools report their absence through check_base_init.
e2e 604 asserts the fresh-home behaviour and, as its control, that a
bootstrapped home announces nothing.

* feat(build): the probe channel, and the device axis as a build program sees it

A rule package is the thing that knows how to ask a machine what it has,
and the engine is the thing that must not. Two directives close that gap:
`mcpp:fact=<name>=<version>` states what the program established about the
machine, `mcpp:floor=<name> >= <version>` what the package needs of it. Both
fold into the runtime declarations a manifest could have carried itself, so
the existing version-floor check reads one list; it now runs a second time
after the root's build program, which is where a rule package speaks from.
Protocol 7.

MCPP_ACCEL carries the resolved device axis (`--accel` / `--no-accel` over
`[build] accel`) to build programs, so an architecture set is written once,
in the manifest. The same value now feeds the `cfg(accelerator = "...")`
layer key, a field that was declared, documented and never written. The
axis also enters the fingerprint and disqualifies the fast path: measured
before, `mcpp build --no-accel` after a device build finished in 0.00 s and
handed back the device build.

* feat(build): a source glob may carry the accel it is for, and the build narrows to it

`[build] sources` accepts `{ glob = "...", accel = "..." }` entries. The glob
joins the plain list, so every reader keeps working; the constraint is
resolved in prepare after feature application. A constrained glob that
matches nothing is refused, naming it. Under a build that asks for no
accelerator the glob is excluded through the same `!` mechanism feature gates
use, which is how one project yields its CPU-only variant. Under a build
that targets something the constraint is not within, the build is refused
naming the glob and both sides (`accel-mismatch`). Device-kind files the
effective set matches reach the build program as MCPP_DEVICE_SOURCES; the
engine has no compile rule for them and the rule package turns each into an
action. Unit tests cover the table form and its refusals; e2e 606 measures
the four outcomes with a backend nothing knows.

* docs(plan): record the engine batch, the fingerprint drift, and the local 168 reading

* feat(build): what a second compiler needs, and the device axis on run and test

Three gaps a rule package that drives a compiler mcpp did not resolve fell
into, each measured on the CUDA example.

**The flags mcpp passes to its own compiler.** `MCPP_TOOLCHAIN_SYSROOT` and
`MCPP_TOOLCHAIN_BINUTILS_DIR` state the `--sysroot` and the `-B` directory,
empty when mcpp passes none. They are not `MCPP_TARGET_SYSROOT`, which is a
tier fact and empty on a hosted target: under a sub-OS the C library is not at
`/usr/include` and the assembler is not at `/usr/bin`, so a compiler that mcpp
did not place fails at the first `#include`:

    crt/host_config.h:218: fatal error: features.h: No such file or directory

`hipcc`, `-fsycl-host-compiler` and any generator that compiles what it emits
have the same gap, so the answer belongs to the engine. Both read from the
single producer that already decides them for the engine's own command lines;
`gcc::binutils_prefix_dir` now states the `-B` guard once, where three copies
of it stood and one said so in a comment.

**A dynamic build-program helper on Linux gets `DT_RPATH`.** The driver's
default is the new tag, and a runpath is consulted only for the helper's own
needed libraries. A build program that opens a host library at run time then
fails one hop later: measured, `dlopen("<sentinel>/lib/libcuda.so.1")`
answered `libdl.so.2: cannot open shared object file` while the very
directories that hold it sat in the helper's RUNPATH. The artifacts mcpp links
already carry DT_RPATH for this reason. The link policy is part of the helper's
cache identity, so an older helper is rebuilt rather than replayed.

**`--accel` / `--no-accel` on `run` and `test`.** The axis existed on `build`
only, so a project could build its CPU-only variant and not run it. Both verbs
take the same two flags, with `--no-accel` travelling as the same explicit
sentinel, and both bypass their fast path when either is given — a cached
artefact was built for whatever axis the last build used.

e2e 607 covers the chained-action shape a device link needs: an artifact-role
action's output consumed by an object-role action, with the intermediate
absent from the link line.

* refactor(doctor): the device toolkit report belongs to the package that owns the tool

`mcpp self doctor` grew a CUDA section: it located a toolkit payload, read
NVIDIA's `crt/host_config.h` for the host-compiler bound, and parsed
`nvcc --dryrun` to name a back-end stage the tool could not reach. Every one of
those readings was correct, and none of them belonged to the engine.

The repository already states the rule as an invariant. `test_runtime_contract`
refuses a vendor name beside a probe launch in `src/`, and the reason is that
an engine which learns to run one vendor's tool learns to run four: the second
backend arrives as a second section, the third as a third, and the engine
acquires a table of tools it must keep current with releases it does not
control.

The same answers are now produced where the tool is known — the rule package
the project imports — and reach mcpp as declarations through the build
program's channel, which the engine compares without knowing what any of the
names mean. Nothing is lost: the CUDA example reports the driver relation, the
host-compiler bound and the unreachable stage before its first compile, and the
version floor refuses a build the machine cannot run.

`test_core_vendor_probes` states the property the removal establishes, over
comment-stripped sources so that recording a vendor's name in a comment stays
possible. It carries its own denominator: an enumeration that found fewer than
a hundred files is a broken scan rather than a clean result.

e2e 602 is removed with the section it tested. Its subject was the doctor
report, not a build, and the same preference — payload before host — is what
the example's rule package now exercises end to end.

* feat(example): two routes for the CUDA example, and the pairings it refuses

The example took one route — nvcc driving the project's own compiler — and
that route has two constraints neither the engine nor the project controls.
Both now produce a sentence before anything is compiled, and a second route
exists that has neither.

**clang is the primary route.** `[toolchain] default = "llvm@22.1.8"` and the
device unit is compiled by the same compiler as the rest of the project
(`-x cuda --cuda-path=<payload>`). No second host compiler, no host-compiler
bound, no CUDA host header in the way. `MCPP_EXAMPLE_CUDA_ROUTE` selects the
other route, and the rule declares `rerun_if_env_changed` for it.

**nvcc is the alternate, and it refuses two pairings by name.** The host
compiler must satisfy the bound nvcc states in its own `crt/host_config.h`: the
rule uses the project's toolchain when it fits, otherwise a `xim:gcc` payload
the project declared, otherwise a refusal naming the declaration to add.
Measured — GCC 16 under nvcc 12.9 fails inside GCC's own `<type_traits>` even
with `-allow-unsupported-compiler`, which admits a compiler one step past the
bound and not a standard library two majors newer.

The second pairing is a toolkit older than the C library. Toolkit 12.9's
`crt/math_functions.h` redeclares the C23 functions `cospi`, `sinpi` and
`rsqrt` for the host without `noexcept`; glibc 2.41 and later declare them with
it, and since C++17 that is part of the function type. The compile stops with
six `exception specification is incompatible` errors naming a glibc header and
a CUDA header, and no decision. The rule reads the C library's
`bits/mathcalls.h` through `mcpp::toolchain_sysroot()` and states the pair it
cannot have, naming the 13.x toolkit as the way out.

**A CPU implementation behind the same seam.** `src/cpu/saxpy.cpp` is selected
by `cfg(not(accelerator = "cuda"))` while the `.cu` carries the accel its glob
is for, so `mcpp build --no-accel` compiles one and `mcpp build` the other,
with no hand-written condition on either side.

Measured on an RTX 4080, driver 550.144.03 reporting CUDA 12.4: `mcpp run` and
`mcpp run --no-accel` both print `12 24 36 48`, from different artifact
directories, and the second contains no `cudaMalloc`. The nvcc route is not
exercisable there — 12.9 meets the driver and not the C library, 13.3 meets the
C library and not the driver — and both refusals are the ones above.

* fix(build): key a package's device sources by its root, not by its bare name

Two packages in one graph may share a bare name and differ only by namespace —
that is what namespaces are for — and the map that carries device sources to a
build program was keyed by the bare name at both its write site and its two
read sites. A collision would hand one package's device sources to the other's
build program, and nothing would report it: the receiving program would emit
actions for files it does not own, and the owner would emit none.

Keyed by the package root instead, which is unique by construction. Found by
re-reading the diff before merge; no test covered it, and a test that did would
need two same-named packages in one graph, which is worth adding when a second
consumer of this map appears.

* release: 2026.9.5.2

The version annotations in the manual said 2026.9.6 while the release lands
today, so they name the version that actually carries them. `mcpp.toml` and
`MCPP_VERSION` move together, as `01_help_and_version.sh` requires.

The CHANGELOG's Unreleased section described the doctor's `nvcc --dryrun`
check as a shipped feature. That check left the engine in this same batch — the
reading was right and the place was wrong — so the entry is rewritten to say
where it went rather than to announce something the release does not contain.

* fix(build): "no accelerator" is the empty string, not the display spelling

`accel_str` prints `(none)` for an empty set so an ABI tag reads as a
sentence. `resolvedAccel` handed that spelling on as a value, and two readers
were wrong at once:

- A build program saw `MCPP_ACCEL=(none)` while the manual promised an empty
  string, so a rule package asking whether there is an accelerator got a yes
  and a backend named `(none)`. This reached every project that never
  mentioned an accelerator, not only those passing `--no-accel`.
- The fingerprint's own guard, `if (!accel.empty())`, was true for every
  project, appending `#accel=(none)` to builds that had asked for nothing —
  the opposite of what the comment beside it says.

Measured with a build program that wrote the value to a file, which is the only
way to see it: mcpp shows a build program's stdout only when it exits non-zero.

e2e 605's fourth section was titled "--no-accel empties both the variable and
the layer" and measured only the layer, which is how this passed a suite it was
meant to be covered by. It now asserts both, and the fixture prints the value
delimited — an empty warning is indistinguishable from an absent one.

* docs(plan): batch ④ landed, and three readings the engine batch produced

The adapters and the rename are in mcpp-index#347; the plan table says so and
records the three findings that came out of measuring them — the display
spelling that leaked into a value, the two upstream couplings the CUDA recipes
now carry, and the path-index namespace behaviour that made local verification
look like a resolution failure.

The verification script travels with the plan rather than living only in a
scratch directory: it is the thing that found the `(none)` defect, and it is
what the release is checked with inside a sandbox.

* refactor(example): the CUDA rule package moves into the mcpplibs namespace

The rule is content this ecosystem wrote rather than anyone's upstream, so
`mcpplibs` is its namespace by the same rule the rest of the index follows, and
the module follows the namespace: `mcpplibs.rules.cuda`.

It also makes the package publishable. The index entry for it points at this
same directory inside an mcpp release tarball — the shape `grpcgen` already
uses — and a descriptor's identity has to match the manifest it points at, so a
package that stays in an `example` namespace can only ever be copied into the
index rather than referenced there.

Measured: `mcpp run` and `mcpp run --no-accel` both print 12 24 36 48 after a
clean rebuild.

* docs(design): §12 — what the implementation overturned, produced, and left

Five places the plan was wrong and two defects the implementation created and
fixed, each with the reading that settled it. Also what a consumer writes after
this batch, and the list of what remains with the reason each item is not done
rather than a silence that reads as done.

* feat(build): an object action's outputs reach a static library, not only a linked image

A package whose device code is its point declares `kind = "lib"`. Until this,
the actions its build program emitted were dropped — with a warning, and with
the archive coming out containing none of them:

    build.mcpp action 'cuda:wkv' has role = "object" but this build produces
    no executable, shared library or test binary to link its outputs into

Measured on llama.cpp's CUDA backend, which is 305 `.cu` files behind exactly
such a target: every action reported that line and the build succeeded, having
produced a library with no device code in it.

The archive rule already consumes `lu.objects`; a static library was simply
absent from the predicate that decides which link units an object action
attaches to. The objects an action produced belong there for the same reason a
compiled `.cpp`'s do — a target's content is what it was told to contain.

e2e 608 asserts the member list of the archive rather than the exit status: an
`ar` handed nothing still writes a well-formed archive and reports success.
This is the engine half of the multi-device design's C-6.

* docs(plan): T5.1 as a gate paid for itself, and what it left

The framework batch's first item found the engine gap C-6 named and nothing
had measured, and it got as far as compiling before stopping on a payload
version rather than on anything about the design. Both readings are recorded
with what would close them.

* fix(example): the clang route passes NVIDIA's own libc++ escape

A device unit that includes <cuda_runtime.h> stops on

    crt/host_defines.h:67: error: "libc++ is not supported on x86 system"

whenever the toolchain is LLVM, which is the toolchain the clang route exists
for. The guard is `#if defined(__CUDACC__) && … && defined(_LIBCPP_VERSION)`,
and clang defines `__CUDACC__` when it compiles CUDA itself, so the refusal —
which is about nvcc's host pass — lands on a compiler it was not written about.

The escape hatch is upstream's own and is passed only on the clang route: nvcc
really does break against libc++, and nothing here weakens that refusal.

Found on ggml's CUDA backend. The example's own kernel never showed it because
a bare kernel includes no toolkit header at all — which is worth knowing about
the example: it exercises the plumbing, not the headers.

* docs(plan): where T5.1's CUDA lane actually stops, and what it proved on the way

The chain the design is about — axis, narrowed glob, device source list, rule
package, actions, static archive, link — runs end to end and produced 48 device
objects. What stops it is a four-way payload matrix that has no satisfiable
point on this machine, and none of its four sides is mcpp's.

* test(e2e): 188's no-consumer case needs a target that does not exist, not a library

A `kind = "lib"` target used to serve as "nothing that takes an object
action's outputs", because a static library did not take them. It does now, so
that fixture stopped exercising the case and the suite reported it — which is
the whole value of running it.

The one shape that still leaves an object action with nowhere to go is a target
behind an inactive `required_features`: it is absent from the build entirely.
A control runs the same project with the feature on and asserts the diagnostic
is gone, so the assertion cannot pass on a build that stopped emitting it for
everyone.

* docs(design): §12.5 — what the framework gate actually measured

T5.1 did what a gate is for: it found the engine gap C-6 named and nothing had
exercised, and it carried the whole chain to 48 device objects before stopping
on a four-way payload matrix with no satisfiable point on this machine. The
matrix is written out with the reading for each side, and the three general
findings it produced are recorded with where they went.

* docs(example): the clang route does pass one flag, and the README should say which

The README claimed "no CUDA host header in the way", which stopped being true
the moment a device unit includes <cuda_runtime.h>. It says what the flag is,
why the guard fires on a compiler it was not written about, and why this
example's own kernel never showed it.

* test(verify): section E asserted a silence that a fresh sandbox gives for free

Rehearsed inside a sandbox against the previous release — the one that still
had the doctor's toolkit reader — and section E reported ok. It was watching
for a section `self doctor` only prints when a CUDA toolkit is installed, and
a fresh sandbox has none, so the assertion could not fail there.

It now reads the engine binary for the names the removed code carried.
`crt/host_config.h` appears twice in 2026.9.5.1 and not at all in 2026.9.5.2,
so the check discriminates: the control run against the older binary fails with
that exact count. The source-level property stays where it belongs, in
`test_core_vendor_probes`.

Also fixed while writing it: `grep -c` prints 0 and exits 1 when it matches
nothing, so `|| echo 0` appended a second line, `[` said "integer expression
expected", and the section reported ok having skipped every comparison.

---------

Co-authored-by: speak-agent <248744407+speak-agent@users.noreply.github.com>
Co-authored-by: speak-agent <x.d2learn.org@gmail.com>

Author: Sunrisepeak

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

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