feat: a CUDA kernel behind a seam module, and the package-level declaration

The example is the design's form A end to end: a device translation unit that
never enters the module graph, an extern "C" interface free of standard-library
types, a seam module that turns that interface back into a C++ one, and a
build-rule package that owns every fact about nvcc's spelling.

The seam is the part worth reading. Its reason for existing is not that nvcc
rejects modules; it is that this is the one place a backend can be exchanged
without any consumer changing, and the one place a cfg(accelerator = ...)
section has somewhere to apply.

The rule package reads the toolkit's own host-compiler bound and selects a
compiler that satisfies it, because mcpp's payload is routinely newer than what
a given toolkit accepts. On the machine this was verified on the payload is gcc
16.1.0 and the toolkit refuses anything above gcc 12, so the rule selects
clang++-14 and says so.

Verified on an RTX 4080 with CUDA 12.0 and driver 550.144.03: mcpp run prints
12 24 36 48, which is the saxpy the kernel computed on the device.

[package] accelerators is declared alongside, mirroring platforms: a statement
of intent and a CI-matrix hint. It is deliberately a different field from an
artifact's accel, because a declaration is written by hand and can be
aspirational while the artifact field is measured from the build and is what a
consumer is refused against.

The example's README records the one hop that is not closed: [xlings] deps
provisions a payload and puts its bin/ on PATH, which is what a payload
providing a program needs, while a payload providing a library also needs its
lib/ on the artifact's runtime search path. Until an mcpp-index package
declares that directory for the driver sentinel, the run needs it supplied.

Author: speak-agent

examples/09-cuda-kernel/README.md

@@ -0,0 +1,98 @@
+# 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.
+
+## The one thing that is not closed yet
+
+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.
+
+xim already has the right shape for this: `libcuda-host-link` is a sentinel
+package that installs a symlink to whatever the host has, so every GPU consumer
+reads one path instead of reimplementing an `ldconfig` probe. This example
+declares it under `[xlings] deps`.
+
+What is missing is the last hop. `[xlings] deps` provisions the payload and
+puts its `bin/` on `PATH`, which is what a payload providing a *program* needs.
+A payload providing a *library* also needs its `lib/` on the artifact's runtime
+search path, and mcpp's private loader does not consult `/usr/lib`, so without
+that hop the statically linked CUDA runtime cannot `dlopen` the driver and
+reports it as "driver version is insufficient". Until an mcpp-index package
+declares that directory, the run needs it supplied:
+
+```
+$ LD_LIBRARY_PATH=$(xlings pkginfo libcuda-host-link)/lib mcpp run
+```
+
+This is an ecosystem gap, not a defect in the design: the durable fix is a
+`compat.cuda-driver` package in mcpp-index that wraps the sentinel and declares
+its library directory, which is the same shape every other runtime provider in
+that index already has.

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<std::string> sources{ "src/kernels/saxpy.cu" };
+    return example::rules::cuda::compile(sources, opt) ? 0 : 1;
+}

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

examples/09-cuda-kernel/app/mcpp.toml

@@ -0,0 +1,42 @@
+[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 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"

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<std::vector<float>>
+saxpy(float a, std::span<const float> x, std::span<const float> y) {
+    if (x.size() != y.size()) return std::nullopt;
+    std::vector<float> out(x.size());
+    if (saxpy_device(a, x.data(), y.data(), out.data(),
+                     static_cast<unsigned>(x.size())) != 0)
+        return std::nullopt;
+    return out;
+}
+
+} // namespace app

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 <cuda_runtime.h>
+#include <cstdio>
+
+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<size_t>(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;
+}

examples/09-cuda-kernel/app/src/main.cpp

@@ -0,0 +1,12 @@
+import std;
+import app.saxpy;
+
+int main() {
+    const std::vector<float> 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<float> want{12, 24, 36, 48};
+    return *out == want ? 0 : 1;
+}

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"