examples: a SYCL kernel and a HIP kernel behind the seam example 09 established (#575)

* examples: a SYCL kernel and a HIP kernel behind the seam example 09 established

Two examples, and the point of both is that they are example 09 one file name
apart. The seam, the CPU fallback, the constrained glob and the rule mechanism
are identical; what differs is which compiler consumes the island.

11-sycl-kernel is the one that tests the claim. Its island is ordinary C++ --
no dialect to see, no launch syntax -- and what makes it a device translation
unit is that it goes to a compiler with a device back end. It also carries the
two properties a SYCL build has and the others do not: TWO EDGES, because a
SYCL object holds its device image and nothing registers it until
`-fsycl-link` emits a wrapper; and TWO C++ RUNTIMES, because `libsycl.so` is
compiled against libstdc++ while an mcpp artifact links libc++. Example 09's
island can promise not to touch the standard library at all; this one cannot,
since SYCL is a C++ library, so the discipline moves from "no standard library"
to "nothing crosses the seam" -- the island catches its own `sycl::exception`
and returns a code.

12-hip-kernel is the same kernel with every device call spelled `hip*`. On the
NVIDIA platform HIP is a header layer over the CUDA runtime, so the compiler is
the project's own clang, `xim:hip-nvidia` contains no binaries, and the
`[dependencies.compat]` entry is the same one hop example 09 needs. Its
`accel` is two chunks -- `hip, cuda12.9+{sm_89}` -- and the second is character
for character what example 09 writes, because a device has one spelling in this
ecosystem however many programming models reach it.

Both answer `12 24 36 48` on an RTX 4080 and again under `--no-accel`, with
zero `/usr` paths on any command line.

The verification script gains three corrections a dry run made, two of them the
same shape -- a criterion pointed at the wrong object:

  * `ldd` prints the program INTERPRETER with `=>` like everything else, so a
    filter looking only for that arrow counted `/lib64/ld-linux-x86-64.so.2` as
    a host library and failed a payload that was in fact complete.
  * A program that cannot START leaves its loader error on the same stderr the
    adapter check reads, so one defect was reported three times and the third
    report named the wrong cause. The adapter checks are skipped when the
    program did not start.
  * A payload may be installed under either namespace, so the lookups accept
    both -- the same two-namespace form xim-pkgindex's own tests use.

It also gains the assertion that would have caught the dpcpp regression: the
programs carry DT_RPATH and the payload's libraries carry no search path at
all, because a RUNPATH on one of those switches off the inherited RPATH of
whatever loaded it.

* examples: 11's island reports what it can, and its README says what it cannot

Synchronised with the fixture the same island came from. Three things were
needed to make "the island returns a code" true, and only two of them are a
catch: the catch inside the buffer scope, the queue's asynchronous handler --
without which SYCL's default handler calls std::terminate, which no catch can
intercept -- and, for the third, a build-time warning from the rule, because a
missing device image is thrown from inside the SYCL scheduler and reaches
neither path.

* docs(plan): 7.7 -- a promise a comment makes about error handling is only tested by the configuration that fails

Two hypotheses were wrong before the backtrace settled it, and both were worth
fixing anyway because each terminates on some machine. The actual cause reaches
neither the caller's frame nor the queue's asynchronous handler, so the rule
states it at build time instead.

* docs(plan): 7.8 -- the ecosystem review, and what each open item would take

Four properties a consumer can rely on across all four lanes, each enforced by
something rather than asserted: the device is spelled once, the engine knows no
vendor name, no device command line reaches the host, and a payload is
reachable rather than merely installed. Three open items with the measurement
that would close each.

* examples: 12 declares the profiler-api payload its device compile needs

Found by the host verification run, not by building on this machine.
`nvidia_hip_runtime_api.h` includes <cuda_profiler_api.h> at its second line
and CUDA ships that header in its own component, so a machine with a host CUDA
installation supplies it from /usr/include and the build works while depending
on something it never declared. The fixture in mcpp-plugins gained the entry
when CI refused it on a runner; the example is the same shape and had the same
gap.

V4 with this: 0 assertions failed -- both examples answer 12 24 36 48 on the
device and again under --no-accel.

* examples: 11 declares one compat entry, and section G reads the farm it built against

Two findings from the sandbox run, and the second is the more useful.

EXAMPLE 11 STILL CARRIED THE PRE-2026.09.07 SHAPE: `compat:cuda-runtime`
beside `compat:sycl-runtime`, plus the driver sentinel. The whole point of the
adapter's 2026.09.07 version is that a SYCL project does not declare CUDA --
the SYCL runtime's back end dlopens the driver, so the adapter carries that
hop. The fixture in mcpp-plugins was updated when the adapter was; the example
was copied from it before that and kept the old form.

SECTION G WAS READING A FARM THE EXAMPLE HAD NOT USED. It picked one out of the
store by `sort -V | tail -1`, and a store that has seen two adapter versions
holds two farms. In the sandbox that reported `ok` for every soname of the
2026.09.06 farm while the example under test built against a different one --
the same shape the round-3 notes record as "the report was chosen by
`ls | head -1`", and the reason those notes exist.

It now reads the directory the example's own `resolution.json` names, and says
so out loud when it cannot: a run with no example build prints `note: falling
back to the newest farm in the store` rather than quietly measuring something
else. Two attempts were needed -- the first looked for `target/*/resolution.json`
where the file is two levels down, which the note caught rather than hid.

A property this exposed and the recipe comment gets wrong is recorded in a
follow-up: `compat.sycl-runtime` has ONE `install()` for both versions and
never reads `pkginfo.version()`, so 2026.09.06 does NOT "install the farm it
always did" -- it installs whatever the current recipe installs. Only the
descriptor entry is frozen.

V4 after both fixes: 0 assertions failed, with section G naming the 2026.09.07
farm and no fallback.

* docs(plan): the round's rows closed with their measurements, and 7.9/7.10

7.3 now carries what each row was decided by rather than that it was done. 7.9
lists the eight defects with the method that found each, because the method is
the reusable part -- and names the shape four of them share: a criterion
pointing at the wrong object, which then answers ok. The mechanical lesson is
that a check which selects its own object must print which one it selected;
section G now does, and that line is what turned the last of these from a false
pass into a visible one.

7.10 records the one red this round did not cause: ci-fresh-install fails in
apt-get inside a debian:11 container on five consecutive commits, two of them
before this round. Diagnosed (the 404'd URL returns 200 from here, so a CDN
node was at fault) and deliberately not patched, because it cannot be verified
from this machine.

---------

Co-authored-by: speak-agent <x.d2learn.org@gmail.com>

Author: Sunrisepeak

.agents/docs/2026-09-05-heterogeneous-build-ecosystem-design-v2.md

@@ -451,8 +451,8 @@ Status is one of `done`, `open`, `deferred (reason)`.
 | N2 | `examples/11-sycl-compute`: one SYCL kernel behind a seam, the CUDA backend and the host device from one artifact | `12 24 36 48` on the RTX 4080 and again with `--no-accel` | N1, Q1 | open |
 | N3 | `examples/12-hip-compute`: a `.hip` kernel through `mcpp.rules.hip` on the NVIDIA platform | `12 24 36 48`, and no `/usr` path on any command line | Q2, Z2 | open |
 | N4 | Chapter 20 gains the SYCL and HIP lanes and the table of which rule drives which compiler; both languages | `check_docs_style.sh` passes; the Chinese chapter has the rows the English one has | - | done: the lane table, the two-chunk `accel` explanation, the HIP header-layer note and the two-C++-runtimes note, in both languages; "Not implemented" corrected -- the whole-target shape is no longer among the missing |
-| N5 | Version 2026.9.6.1, CHANGELOG, unit and e2e suites, PR, CI green, self-review | no e2e failure other than the ones already failing on `origin/main` | N1-N4 | open |
-| N6 | Release, mirror, index bump, bootstrap pin | both mirrors return 200 and identical bytes for every asset; xim-pkgindex names it `latest` | N5 | open |
+| N5 | Version 2026.9.6.1, CHANGELOG, unit and e2e suites, PR, CI green, self-review | no e2e failure other than the ones already failing on `origin/main` | N1-N4 | done: PR #574, 36/36 checks green, merged as `12d4759`; main green on every workflow at that sha except `ci-fresh-install`, which fails in `apt-get` inside a debian:11 container on five consecutive commits including two that predate this round |
+| N6 | Release, mirror, index bump, bootstrap pin | both mirrors return 200 and identical bytes for every asset; xim-pkgindex names it `latest` | N5 | done: all four artefacts byte-identical GitHub/GitCode by GET; bump PR #770 merged (`6c2f283`) with `Publish Index Artifact` green; `xlings install mcpp@2026.9.6.1` yields `mcpp 2026.9.6.1` at the store path; pin on main |
 
 #### mcpp-plugins (single PR, version 0.2.0)
 
@@ -461,22 +461,22 @@ Status is one of `done`, `open`, `deferred (reason)`.
 | Q1 | `mcpp.rules.sycl`, feature `rules-sycl` | a fixture compiles a `.sycl` unit with the `dpcpp` payload and links it into a C++23-modules program; the rule refuses with the declaration to add when the payload is absent | N1 | done: `12 24 36 48` on an RTX 4080 through `mcpp run`, zero `/usr` paths on any command line, and the device-link wrapper asserted as its own file |
 | Q2 | `mcpp.rules.hip`, feature `rules-hip` | a fixture compiles a `.hip` unit on the NVIDIA platform; the rule names the platform it selected and why | Z2 | done: `12 24 36 48` on an RTX 4080, zero `/usr` paths; the AMD platform is refused by name with the reason rather than approximated |
 | Q3 | `mcpp.rules.spirv` gains the `glslc` route now that a payload exists | both compilers produce a header the same program includes; the rule states which one it used | Z1 | done: the two command lines share almost nothing -- glslc's `-mfmt=c` is a bare initialiser list and its `-S` means "emit assembly" where glslang's names a stage -- so the rule writes the declaration itself and keys its `mcpp::fact` on the flavour |
-| Q4 | CI: one consumer fixture per feature, built with the pinned mcpp | green on the PR head | Q1-Q3, N6 | open |
-| Q5 | Release `v0.2.0`; GitHub archive and a GitCode asset with identical bytes; index descriptor | both URLs return 200 and one sha256 | Q4 | open |
+| Q4 | CI: one consumer fixture per feature, built with the pinned mcpp | green on the PR head | Q1-Q3, N6 | done, and it found three host leaks the previous check could not see -- see 7.9 |
+| Q5 | Release `v0.2.0`; GitHub archive and a GitCode asset with identical bytes; index descriptor | both URLs return 200 and one sha256 | Q4 | done: PR #3 merged (`56d9da2`), tag v0.2.0, both mirrors 47,792 bytes under one sha256 (`b5ee0cf4`) |
 
 #### mcpp-index (single PR)
 
 | # | task | criterion | depends on | status |
 |---|---|---|---|---|
 | R1 | `compat.sycl-runtime` 2026.09.06; then `mcpp:plugins` 0.2.0 in `pkgs/m/mcpp.plugins.lua`, floor 2026.9.6.1 | the adapter: a workspace member loads `libsycl.so.9` by soname under mcpp's own loader. the descriptor: `mcpp add mcpp:plugins` resolves in a sandbox and the new features are selectable | Q5, N6 | adapter done (PR #354 merged as `e17fc88`, index artifact published); descriptor open |
-| R2 | CI green; merge; index artifact published | `Publish Index Artifact` green on the merge commit | R1 | open |
+| R2 | CI green; merge; index artifact published | `Publish Index Artifact` green on the merge commit | R1 | done: PR #356 merged (`e68c9d8`), artifact green. The snapshot needed the whole `data/<index>` directory removed to move -- clearing the version marker and the cache JSON was not enough, and neither mirror was at fault |
 
 #### Verification
 
 | # | task | criterion | depends on | status |
 |---|---|---|---|---|
-| V3 | A fresh subos, CN mirror for both tools: install mcpp 2026.9.6.1 from the index, build examples 09 through 12 against `mcpp:plugins` 0.2.0 | every example builds; example 10 answers on lavapipe; the device half of 09, 11 and 12 compiles and says cleanly that the sandbox has no device | Z4, N6, Q5, R2 | open |
-| V4 | The same on this host with the GPU | `12 24 36 48` from 09, 11 and 12 on the RTX 4080 | V3 | open |
+| V3 | A fresh subos, CN mirror for both tools: install mcpp 2026.9.6.1 from the index, build examples 11 and 12 against `mcpp:plugins` 0.2.0 | every example builds; the device half compiles and the sandbox says cleanly that it has no device | Z4, N6, Q5, R2 | **done, 0 assertions failed.** Both examples built and answered `12 24 36 48` under `--no-accel`; the device-link wrapper was produced; no `/usr` on any command line. It also found two defects of mine -- see 7.9 |
+| V4 | The same on this host with the GPU | `12 24 36 48` from 11 and 12 on the RTX 4080 | V3 | **done, 0 assertions failed.** Both examples on the device, `sycl-ls` enumerating the RTX 4080 from the store, and the adapter farm complete down to `libcuda.so.1` |
 
 ### 7.4 The completion rule
 
@@ -587,3 +587,167 @@ mcpp-plugins, then mcpp-index (descriptor), then mcpp (examples).
 | 2026-09-06 | the round is five PRs, not three | the release order is a cycle otherwise |
 | 2026-09-06 | dpcpp corrected a second time (xim #769) | the first correction fixed the payload's own programs and broke every consumer of it; the regression was invisible from the payload |
 | 2026-09-06 | `compat.sycl-runtime` 2026.09.07 (mcpp-index #355) | a SYCL project should not have to declare CUDA; the adapter owns the hop its runtime needs |
+| 2026-09-06 | the SYCL island gains an async handler and an inner catch; the rule warns when `accel` names sycl and no device | exercising the rule's other branch showed the island could not keep the promise its own comment made; see 7.7 |
+
+### 7.7 What exercising the second branch found
+
+`mcpp.rules.sycl` has two branches: `accel = "sycl, cuda12.9+{sm_89}"` compiles
+ahead of time for that architecture, and `accel = "sycl"` compiles to SPIR-V and
+leaves the device to the runtime. The examples take the first. Taking the second
+built correctly and then died on an RTX 4080 with
+
+    terminate called after throwing an instance of 'sycl::_V1::exception'
+    terminate called recursively
+
+and no message of the program's own -- contradicting the fixture's own comment,
+which said the island turns its failures into a return code.
+
+Two hypotheses were wrong before the backtrace settled it, and both were worth
+fixing anyway because each terminates on some machine:
+
+* An exception unwinding through the three `sycl::buffer` destructors. A buffer
+  destructor blocks until the work that reads it has finished, so a catch
+  outside the scope unwinds through three destructors that are themselves
+  finishing failed work. The catch moved inside the scope.
+* The queue's asynchronous handler. SYCL delivers most device errors to it, and
+  a queue constructed without one gets the default handler, which calls
+  `std::terminate`. No catch intercepts that, because it never travels as an
+  exception through the frame. A handler is installed.
+
+The actual cause is neither. `ProgramManager::getDeviceImage` throws inside the
+SYCL scheduler because the CUDA back end does not consume SPIR-V, and that path
+reaches neither the caller's frame nor the queue's handler. Nothing the island
+can write catches it.
+
+So the rule says it at build time instead, as an advisory rather than a
+refusal: the SPIR-V form is correct and portable, and the rule cannot know the
+machine's devices -- which is exactly why the statement belongs before the
+first kernel rather than in a crash with no message.
+
+The general shape, which is the reusable part: **a promise a comment makes
+about error handling is only tested by the configuration that fails.** Every
+run on this machine took the path where nothing goes wrong, and the comment
+read as true for as long as that held.
+
+### 7.8 Ecosystem review: what a consumer sees after this round
+
+A project that wants a GPU now writes the same four things whichever
+programming model it picks, and the fourth is the only one that names the
+model.
+
+```toml
+[dependencies.mcpp]
+plugins = { version = "0.2.0", features = ["rules-sycl"], host-module = true }
+
+[dependencies.compat]
+sycl-runtime = "2026.09.07"     # the runtime the artifact reaches through the loader
+
+[xlings.workspace]              # payloads: the project picks the versions
+"xim:dpcpp"     = "7.1.0"
+"xim:gcc"       = "15.1.0"
+"xim:cuda-nvcc" = "12.9.86"
+
+[build]
+accel = "sycl, cuda12.9+{sm_89}"
+```
+
+Four properties hold across all four lanes, and each is enforced by something
+rather than asserted here.
+
+**The device is spelled once.** `cuda12.9+{sm_89}` is character for character
+the same in the CUDA, HIP and SYCL manifests. The first chunk names the
+programming model; the second names the device. A rule reads the chunk it owns
+and ignores the rest, so adding a fifth model adds no spelling.
+
+**The engine still knows no vendor.** `test_core_vendor_probes` strips comments
+from `src/` and refuses any vendor tool name, with the file count as its own
+denominator. Three rules and four payloads later, that test is unchanged.
+
+**No device command line reaches the host.** Asserted per lane in the plugin
+collection's CI as a separate step from the build, because all three lanes
+compiled successfully on a developer machine while reading `/usr/include` --
+twice in this round alone, once for the SYCL unit's C++ standard library and
+once for clang's CUDA installation. Measured here: zero `/usr` paths across
+CUDA, HIP and SYCL.
+
+**A payload is reachable, not merely installed.** The round's largest single
+defect was a payload whose five reporting programs could not start; the round's
+second was the correction to it, which fixed the payload and broke every
+consumer. Both are now assertions: the programs start and carry DT_RPATH, and
+the payload's libraries carry no search path at all.
+
+Three things this round could not close, each with what would decide it:
+
+| open | what would decide it |
+|---|---|
+| Intel `anv` in the Mesa payload | a libclc and SPIRV-LLVM-Translator chain, then a Mesa rebuild; both packages exist on conda-forge and were confirmed repackable, and are withheld because nothing yet consumes them |
+| the AMD platform of HIP | a ROCm runtime and device library in `xim-pkgindex`. The rule refuses it by name today rather than emitting an object nothing can link |
+| llama.cpp's Vulkan lane | its blocker is gone -- `glslc` is published -- and the remaining work is in that project's own `-m` repository rather than here |
+
+And one thing measured rather than assumed: `accel = "sycl"` without a device
+compiles to SPIR-V, which the CUDA back end does not consume. The build is
+correct, the failure is the machine's, and the rule now says so before the
+first kernel. See 7.7.
+
+### 7.9 What CI and the two verification runs found, and the shape they share
+
+Eight defects in this round, six of them in work written for it, and not one
+was found by reading the code. The list is worth keeping because the *method*
+that found each is the reusable part.
+
+| found by | defect |
+|---|---|
+| using the payload | `xim:dpcpp` shipped five programs that could not start |
+| CI's dependency-closure check | the first repair swapped the interpreter, losing the host fallback |
+| rebuilding the example | the second repair fixed the payload and broke every consumer of it |
+| reading `clang -###` | `-lstdc++` is rewritten to `-lc++` and vanishes |
+| a runner | `libz.so.1` missing from the farm; this machine had it |
+| taking the rule's other branch | the SYCL island's error promise held only where nothing failed |
+| a runner | the HIP device pass read the host's `/usr/include/c++/14` |
+| the sandbox run | example 11 carried a superseded manifest shape, and section G measured a farm the example had never used |
+
+THE SHAPE MOST OF THEM SHARE is that a criterion pointed at the wrong object,
+and the wrong object answered `ok`:
+
+* the host-leak check read the command line, and an implicit include search is
+  never on it -- so it reported clean on the machine that was leaking;
+* the farm check picked a directory out of the store by sorting, and a store
+  that has seen two adapter versions holds two farms;
+* a payload lookup named one namespace, and the payload was installed under the
+  other;
+* a program that could not start had its loader error attributed to an adapter,
+  turning one defect into three reports of which the third named the wrong
+  cause.
+
+Round 3's notes already record this shape twice ("the report was chosen by
+`ls | head -1`", "the mirror was read from the tool's stdout"). It recurred
+four times here, which suggests the lesson is not "watch for this" but
+something mechanical: **when a check reports on an object it selected, it must
+print which object it selected.** Section G now does, and the line
+`note: falling back to the newest farm in the store` is what turned the last of
+these from a false pass into a visible one.
+
+A property the same run exposed in a merged package: `compat.sycl-runtime` has
+one `install()` and never reads `pkginfo.version()`, so a version key selects
+the anchor URL and nothing else. The 2026.09.07 commit said 2026.09.06
+"installs the farm it always did"; it does not, and mcpp-index #357 corrects
+the comment rather than the behaviour, which is intended for an adapter whose
+content is generated.
+
+### 7.10 One failure this round did not cause and did not fix
+
+`ci-fresh-install` is red on mcpp's main and has been for five consecutive
+commits, two of them before this round began. It fails in `apt-get install`
+inside a `debian:11` container, in a step named "Install prerequisites" that
+runs before mcpp is involved:
+
+    E: Failed to fetch .../libperl5.32_5.32.1-4+deb11u5_amd64.deb  404  Not
+    Found [IP: 146.75.38.132]
+
+Measured rather than assumed: that exact URL returns 200 from here, and
+`bullseye-security/InRelease` is still served -- so the package exists and a
+CDN node was answering 404. The step already runs `apt-get update`, so the
+mitigation would be `-o Acquire::Retries=3` rather than a fresh index. It is
+not applied here: Docker Hub is unreachable from this machine, the failure does
+not reproduce from this network, and an unverifiable change to an unrelated
+workflow is worth less than a precise diagnosis of it.