0.2.0: mcpp.rules.sycl, mcpp.rules.hip, and the glslc route the payload made real (#3)

* 0.2.0: mcpp.rules.sycl, mcpp.rules.hip, and the glslc route the payload made real

Three members, and each is the second instance of something the collection had
built once. A single lane proves a rule package can drive one vendor's
compiler; the claim this ecosystem makes is that the engine knows no vendor at
all, and only a second instance tests it.

mcpp.rules.hip -- THE SECOND DEVICE API.

On the NVIDIA platform HIP is a header layer over the CUDA runtime: every entry
point is an inline wrapper over the CUDA one, so the object links against the
CUDA runtime and there is no ROCm on the machine. The rule therefore drives the
SAME compiler `mcpp.rules.cuda` uses on its clang route -- the project's own
toolchain clang, `-x cuda` -- plus `xim:hip-nvidia`'s include directory and
`-D__HIP_PLATFORM_NVIDIA__`. hipcc is deliberately not used: it is a driver
that reads HIP_PLATFORM and picks nvcc or amdclang, and every decision it makes
is one this rule has already made from the declaration. The AMD platform is
refused by name rather than approximated, because this ecosystem publishes no
ROCm runtime and the object would link against nothing.

mcpp.rules.sycl -- THE SECOND COMPILATION MODEL.

A SYCL translation unit is ordinary C++ and reaches this rule through the same
constrained glob that carries a `.cu`, as `.sycl` (mcpp 2026.9.6.1). It submits
N compile actions and one that consumes their outputs: a SYCL object carries
its device image but nothing registers it, and `-fsycl-link` is what produces
the registration. Three flags exist only to keep the host out, and two of them
were added because a build that already worked was found to be reading it:

  --gcc-install-dir  without it the unit compiled against /usr/include/c++
  --cuda-path        without it clang found the host's CUDA installation
  -x c++             `.sycl` is this ecosystem's spelling and no compiler knows
                     it; without this the driver classifies the file as a
                     LINKER INPUT, warns `'linker' input unused`, exits 0 and
                     produces nothing

Neither of the first two says anything when missing. They are visible only in
the compiler's own include search list, and only on a machine that has those
directories -- which is why the fixture CI step that greps a verbose build for
`/usr` is separate from the step that builds it.

The link carries `-l:libstdc++.so.6` rather than `-lstdc++`, and that is not a
style choice: clang's driver treats `-lstdc++` as a SELECTOR for the C++
standard library and rewrites it to `-lc++` under `-stdlib=libc++`, so the flag
disappears from the link line with no diagnostic and `std::cerr` comes back
undefined from inside a SYCL header. `-l:<soname>` names a file and is not
rewritten.

mcpp.rules.spirv -- THE ROUTE THAT WAS A CLAIM.

This file said glslc was unsupported "because nothing in this ecosystem
publishes it, and a route with no payload behind it is a claim rather than a
feature". `xim:shaderc` publishes it, so the route exists. The two compilers
are not interchangeable: glslang's `-x --vn` emits a complete C declaration
where glslc's `-mfmt=c` emits a bare initialiser list, and glslc's `-S` means
"emit assembly" where glslang's names a stage. The rule writes the declaration
around glslc's output itself, uses `-fshader-stage=`, and records WHICH
compiler it found as a fact keyed on the flavour -- two compilers that produce
an equivalent header from one shader are still two different answers to "what
compiled this".

Measured on an RTX 4080, each through `mcpp run` and each with zero `/usr`
paths on any command line: the HIP fixture answers `12 24 36 48`, and so does
the SYCL fixture. The CPU variant of both answers the same through the same
seam under `--no-accel`.

A SYCL build carries two C++ runtimes -- `libsycl.so` is compiled against
libstdc++ and an mcpp artifact links libc++ -- and mcpp's duplicate-symbol
check reports the unwinder symbols they share. That warning is correct, and it
is why the seam discipline is not optional here: the fixture catches its SYCL
exception inside the device translation unit and returns a code, because the
runtime that threw it is not the one the caller would unwind with.

* ci: the glslc route through the same fixture, and the SYCL fixture's one compat entry

The glslc step is the assertion Q3 was missing: one source, two compilers, one
answer. It builds tests/spirv-consumer a second time with MCPP_GLSLC set,
asserts the `.inc` the glslc route produces (glslang's route produces no such
file, so its presence is what distinguishes the two), and asserts the program
prints the same magic number. Measured locally: 370 words through glslang and
293 through glslc, one program.

The SYCL fixture now declares `compat:sycl-runtime = "2026.09.07"` and nothing
else from compat: the adapter carries the driver hop its own runtime needs, so
a project that writes SYCL does not declare CUDA.

* fix: two diagnostics that told the reader to solve a problem they had solved

Both found by re-reading the new rules rather than by a failing build, and both
are the same shape: a message that is correct in isolation and wrong given what
the caller already did.

`mcpp.rules.sycl` asked a project that had set `options::compiler` to declare
`xim:dpcpp` anyway. The payload is needed for the COMPILER; a project that
named its own does not need it.

`mcpp.rules.spirv` followed its precise refusal of an unrecognised
`options::compiler` -- naming the program and saying that glslang and glslc
share almost no flags -- with the generic "no shader compiler found. Install
one into the workspace", which contradicts it. Discovery now records that it
has already reported, and the caller stays quiet. Measured: one message where
there were two.

Both routes re-verified through tests/spirv-consumer: 370 words through
glslang, 293 through glslc, one program and one magic number.

* ci: 90 minutes, because the SYCL fixture pulls three payloads this job never needed

dpcpp is 578 MB installed, gcc 265 MB and cuda-nvcc 319 MB, and the cache key
includes the fixtures' manifests -- so the run that adds a fixture is always
the cold one. 60 minutes was chosen when the heaviest consumer was a shader.

* fix(sycl): the island reports what it can, and the rule warns about what it cannot

Found by exercising the rule's OTHER branch -- `accel = "sycl"` with no device
chunk, which compiles to SPIR-V and leaves the choice to the runtime. It builds
correctly. It then dies 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, which contradicted this fixture's own
claim that the island turns its failures into a return code.

Two of the three guesses on the way there were wrong, and the backtrace settled
it. It is not an exception unwinding through the buffer destructors, and it is
not the queue's default asynchronous handler: the throw comes from
`ProgramManager::getDeviceImage` inside the SYCL scheduler, because the CUDA
back end does not consume SPIR-V. It travels through neither the caller's frame
nor the queue's handler, so no amount of care in the island catches it.

Both of the wrong guesses were nonetheless real gaps and are fixed, because
each is catchable and each would have terminated on some machine:

  * the catch is inside the buffer scope -- a buffer destructor blocks until
    the work that reads it finishes, and unwinding through three of those is a
    second throw during unwinding;
  * the queue takes an asynchronous handler -- a queue constructed without one
    gets the default, and the default calls `std::terminate`, which is not
    something a catch can intercept because it does not travel as an exception
    through the frame.

For the third, the rule now says so at build time. 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 precisely why it is worth saying before the first
kernel rather than leaving it to a crash without a message.

The example's own manifest names the device, so it is unaffected: both variants
still answer `12 24 36 48`.

* fix(hip): the toolchain's clang by path, its C++ library named, and the header CUDA ships separately

CI found two host leaks that the check meant to prevent them could not see,
and the reason is the same for both: an IMPLICIT include search never appears
on a command line, so grepping a verbose build for `/usr` reported nothing
wrong on a developer machine AND on a runner while the runner was reading the
host's headers.

THE C++ STANDARD LIBRARY WAS THE COMPILER'S DEFAULT. A bare CUDA kernel
includes no standard library header, which is why `mcpp.rules.cuda` never had
to say anything; the HIP headers reach `<limits>`, so the device pass compiles
one and something has to decide which. Measured: clang's own configuration
supplied this ecosystem's libc++ here, and on a runner the same clang fell back
to detecting the host's GCC and read `/usr/include/c++/14`, whose `<limits>`
declares `__float128` -- which the NVPTX device target does not support.
Nineteen errors from a header no part of this ecosystem chose. The rule now
passes `-nostdinc++` and the toolchain's own libc++, finding the per-triple
overlay rather than constructing its name (`x86_64-unknown-linux-gnu`, where
`mcpp::target()` says `x86_64-linux-gnu`).

`cuda_profiler_api.h` WAS COMING FROM A HOST CUDA INSTALLATION.
`nvidia_hip_runtime_api.h` includes it at its second line and CUDA ships it in
its own component. `xim:cuda-profiler-api` is now among the payloads the rule
requires and names.

AND THE COMPILER WAS AN IDENTITY, NOT A PATH. `mcpp::compiler()` answers with
the string `clang`; putting that where an argv[0] belongs runs whichever
`clang` the action's PATH offers. It works on a machine whose PATH already has
the payload and is a host leak everywhere else. The rule takes
`mcpp::toolchain_dir() + "/bin/clang++"`, which is what `mcpp.rules.cuda`
does and for this reason.

The CI step is replaced rather than repaired: it now reads the compiler's own
`#include <...> search starts here` list. Its criterion is the C++ standard
library specifically and not the absence of `/usr` -- mcpp's own compiles leave
`/usr/include` on that list as a last resort for C headers, and a stricter rule
here would be stricter than the engine it checks.

Verified on all three lanes: no host C++ standard library on any of them, and
the HIP fixture still answers `12 24 36 48`.

* ci: remove the superseded host-leak step, which the replacement left behind

The previous commit added the step that reads the compiler's include search
list and was meant to replace the one that greps the command line. Its
end-boundary was wrong, so both shipped and the run carried two assertions
about the same property -- one of which had already been shown unable to see
the defect it existed to catch.

---------

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

Author: Sunrisepeak

.github/workflows/ci.yml

@@ -9,13 +9,18 @@ on:
 env:
   # The mcpp release the consumers build with. Raising it is what admits a
   # member that relies on a newer engine; the README states each member's floor.
-  MCPP_VERSION: 2026.9.5.4
+  MCPP_VERSION: 2026.9.6.1
 
 jobs:
   consumers:
     name: consumers (linux x86_64)
     runs-on: ubuntu-24.04
-    timeout-minutes: 60
+    # 90 rather than 60: the SYCL fixture pulls three payloads this job did not
+    # need before -- dpcpp (578 MB installed), gcc (265 MB) and cuda-nvcc
+    # (319 MB) -- and a cold cache downloads all of them before the first
+    # compile. The cache key includes the fixtures' manifests, so a run that
+    # changes one of them is always the cold case.
+    timeout-minutes: 90
     steps:
       - uses: actions/checkout@v4
 
@@ -47,6 +52,25 @@ jobs:
           "$MCPP" run | tee run.log
           grep -q '^magic=07230203' run.log
 
+      # THE SECOND COMPILER, THROUGH THE SAME FIXTURE. glslang and glslc share
+      # almost no flags -- glslc's `-mfmt=c` is a bare initialiser list where
+      # glslang's `-x --vn` is a complete declaration, and glslc's `-S` means
+      # "emit assembly" where glslang's names a stage -- so the rule writes the
+      # declaration itself on this route. The program is unchanged, which is the
+      # assertion: one source, two compilers, one answer.
+      - name: rules-spirv through glslc
+        working-directory: tests/spirv-consumer
+        run: |
+          curl -L -fsS --retry 3 --retry-all-errors -o shaderc.tar.gz \
+            "https://github.com/xlings-res/shaderc/releases/download/2026.3/shaderc-2026.3-linux-x86_64.tar.gz"
+          tar -xzf shaderc.tar.gz
+          rm -rf target
+          MCPP_GLSLC="$PWD/shaderc-2026.3/bin/glslc" "$MCPP" build
+          test -f target/.build-mcpp/out/spirv/scale_comp.inc
+          MCPP_GLSLC="$PWD/shaderc-2026.3/bin/glslc" "$MCPP" run | tee run-glslc.log
+          grep -q '^magic=07230203' run-glslc.log
+          rm -rf shaderc.tar.gz shaderc-2026.3
+
       # A tool, not a rule: the header is written while the build program runs,
       # so there is no action to schedule. The second build is the measurement
       # that matters -- editing the data file must reach the binary, which is
@@ -72,3 +96,60 @@ jobs:
           "$MCPP" build --no-accel
           "$MCPP" run --no-accel | tee run.log
           grep -q '^12 24 36 48' run.log
+
+      # The same shape one API over: HIP on the NVIDIA platform is a header
+      # layer over the CUDA runtime, so this compiles the same device code
+      # through a different spelling and needs no ROCm. The device object is
+      # asserted to exist, because a rule that submitted no action at all would
+      # otherwise leave a build that "succeeded" and produced nothing.
+      - name: rules-hip through a consumer
+        working-directory: tests/hip-consumer
+        run: |
+          "$MCPP" build
+          test -f target/.build-mcpp/out/saxpy.hip.o
+          "$MCPP" build --no-accel
+          "$MCPP" run --no-accel | tee run.log
+          grep -q '^12 24 36 48' run.log
+
+      # SYCL is compiled by a second compiler, so this step is the one that
+      # would catch a payload the rule cannot drive. Two objects have to exist:
+      # the unit's own, and the device-link wrapper that registers its image --
+      # without the second the program links and finds no kernel at run time.
+      - name: rules-sycl through a consumer
+        working-directory: tests/sycl-consumer
+        run: |
+          "$MCPP" build
+          test -f target/.build-mcpp/out/saxpy.sycl.o
+          test -f target/.build-mcpp/out/sycl_device_link.o
+          "$MCPP" build --no-accel
+          "$MCPP" run --no-accel | tee run.log
+          grep -q '^12 24 36 48' run.log
+
+      # THE HOST-LEAK ASSERTION, and what it measures was corrected once.
+      #
+      # Its first form grepped a verbose build for `/usr`. That reported
+      # nothing wrong on a developer machine AND on a runner, while the runner
+      # was in fact compiling the HIP device pass against the host's
+      # `/usr/include/c++/14` -- because an IMPLICIT include search never
+      # appears on a command line. The compiler's own `#include <...> search
+      # starts here` list is where it is visible.
+      #
+      # The criterion is the C++ standard library specifically, not `/usr`
+      # absence: mcpp's own compiles leave `/usr/include` on the list as a last
+      # resort for C headers, and a stricter rule here would be stricter than
+      # the engine it is checking. What must never appear is a C++ library
+      # nothing in this ecosystem chose.
+      - name: no device compile reads a host C++ standard library
+        run: |
+          fail=0
+          for d in tests/cuda-consumer tests/hip-consumer tests/sycl-consumer; do
+            ( cd "$d" && rm -rf target && "$MCPP" build -v > build.log 2>&1 ) || {
+              echo "build failed in $d"; tail -20 "$d/build.log"; exit 1; }
+            hits=$(grep -oE '/(usr|opt)/[^ ]*(c\+\+|include/c\+\+)[^ ]*' "$d/build.log" | sort -u | head -5)
+            if [ -n "$hits" ]; then
+              echo "host C++ standard library reached from $d:"; echo "$hits"; fail=1
+            else
+              echo "ok: $d reads no host C++ standard library"
+            fi
+          done
+          [ "$fail" -eq 0 ]

README.md

@@ -6,7 +6,7 @@ imports each one from `build.mcpp` under the module name the member declares.
 
 ```toml
 [dependencies.mcpp]
-plugins = { version = "0.1.1", features = ["rules-spirv"], host-module = true }
+plugins = { version = "0.2.0", features = ["rules-spirv"], host-module = true }
 ```
 
 ```cpp
@@ -39,14 +39,19 @@ engine's own module family and is not used here.
 | feature | module | since mcpp | what it needs |
 |---|---|---|---|
 | `rules-cuda` | `mcpp.rules.cuda` | 2026.9.5.2 | the toolkit named in `[xlings.workspace]` (`xim:cuda-nvcc`, `xim:cuda-cudart`, and `xim:libcurand` for the clang route, whose wrapper includes a cuRAND header unconditionally), `[build] accel = "cuda…"`, a constrained glob for `*.cu`; the clang route with an LLVM toolchain, the nvcc route with a GCC one |
-| `rules-spirv` | `mcpp.rules.spirv` | 2026.9.5.3 | `xim:glslang` in `[xlings.workspace]`, `[build] accel = "vulkan1.2"`, a constrained glob for the shader stages; emits one header per shader through a `role = "source"` action |
+| `rules-hip` | `mcpp.rules.hip` | 2026.9.5.2 | `xim:hip-nvidia` plus the CUDA back end it compiles through (`xim:cuda-nvcc`, `xim:cuda-cudart`, `xim:libcurand`, `xim:cuda-cccl`), `[build] accel = "hip, cuda12.9+{sm_89}"`, a constrained glob for `*.hip`. On the NVIDIA platform HIP is a header layer over the CUDA runtime, so the compiler is the project's own clang and there is no ROCm on the machine |
+| `rules-spirv` | `mcpp.rules.spirv` | 2026.9.5.3 | `xim:glslang` or `xim:shaderc` in `[xlings.workspace]`, `[build] accel = "vulkan1.2"`, a constrained glob for the shader stages; emits one header per shader through a `role = "source"` action, and states which of the two compilers produced it |
+| `rules-sycl` | `mcpp.rules.sycl` | 2026.9.6.1 | `xim:dpcpp` (the compiler), `xim:gcc` (the C++ standard library the unit compiles against, not a second toolchain) and `xim:cuda-nvcc` for an NVIDIA target; `[build] accel = "sycl"` or `"sycl, cuda12.9+{sm_89}"`, a constrained glob for `*.sycl`, and `compat:sycl-runtime` so the artifact can reach `libsycl.so.9` at run time. The floor is the release whose device-source table carries `.sycl` |
 | `tools-embed` | `mcpp.tools.embed` | 2026.9.5.4 | nothing beyond mcpp: it reads a file and writes a header while the build program runs. The floor is the release whose fast path compares a declared file input, without which an edit to the data does not reach the binary |
 
 The floor is the mcpp release whose engine carries what the member relies on:
 `rules-spirv` needs the device-source table that classifies shader extensions,
-which 2026.9.5.3 introduced, and `tools-embed` needs the fast path to compare a
-declared file input, which 2026.9.5.4 introduced. The index descriptor states the floor; a project
-on an older mcpp is refused at resolution rather than at the first shader.
+which 2026.9.5.3 introduced; `tools-embed` needs the fast path to compare a
+declared file input, which 2026.9.5.4 introduced; and `rules-sycl` needs `.sycl`
+in that same device-source table, which 2026.9.6.1 introduced. The index
+descriptor states the highest floor among the members, so it is the floor of the
+collection rather than of any one feature; a project on an older mcpp is refused
+at resolution rather than at the first shader.
 
 ## How the engine sees this package
 
@@ -67,6 +72,20 @@ tools/<x>.cppm     export module mcpp.tools.<x>;
 tests/<consumer>/  one project per member, built by CI with the pinned mcpp
 ```
 
+## What a rule may drive, and what it may not
+
+A member drives a compiler the ecosystem resolved and no other. `xim:dpcpp` is
+`mcpp.rules.sycl`'s compiler; `xim:gcc` is the C++ standard library it compiles
+against; `xim:cuda-nvcc` is the back end it emits for. None of the three is a
+default: each is read from `mcpp::xpkg_dir`, and a member that cannot find one
+refuses and prints the `[xlings.workspace]` line that would add it.
+
+Two of those three were added because a build that already worked was found to
+be reading the host. Without `--gcc-install-dir` the SYCL unit compiled against
+`/usr/include/c++`; without `--cuda-path` clang found the host's CUDA
+installation. Neither said anything: both are visible only in the compiler's
+own include search list, and only on a machine that has those directories.
+
 ## Adding a member
 
 1. One file, `rules/<x>.cppm` or `tools/<x>.cppm`, declaring its module name.

mcpp.toml

@@ -1,7 +1,7 @@
 [package]
 name        = "plugins"
 namespace   = "mcpp"
-version     = "0.1.1"
+version     = "0.2.0"
 description = "Official mcpp build plugins: rule packages under mcpp.rules.*, build-time utilities under mcpp.tools.*, each member selected by a feature"
 license     = "Apache-2.0"
 authors     = ["mcpp-community"]
@@ -22,8 +22,10 @@ sources = ["src/plugins.cppm"]
 [features]
 default     = []
 rules-cuda  = { sources = ["rules/cuda.cppm"] }
+rules-hip   = { sources = ["rules/hip.cppm"] }
 rules-spirv = { sources = ["rules/spirv.cppm"] }
-tools-embed  = { sources = ["tools/embed.cppm"] }
+rules-sycl  = { sources = ["rules/sycl.cppm"] }
+tools-embed = { sources = ["tools/embed.cppm"] }
 
 [targets.plugins]
 kind = "lib"