From 56bafa0f844aa3836f4077bf4f54ffdbf5a32261 Mon Sep 17 00:00:00 2001 From: speak-agent Date: Sat, 5 Sep 2026 19:01:33 +0800 Subject: [PATCH 1/3] feat: mcpp:plugins 0.1.0 with mcpp.rules.cuda and mcpp.rules.spirv --- .github/workflows/ci.yml | 59 +++ README.md | 86 +++ mcpp.toml | 28 + rules/cuda.cppm | 619 ++++++++++++++++++++++ rules/spirv.cppm | 356 +++++++++++++ src/plugins.cppm | 16 + tests/cuda-consumer/build.mcpp | 9 + tests/cuda-consumer/include/saxpy/saxpy.h | 21 + tests/cuda-consumer/mcpp.toml | 73 +++ tests/cuda-consumer/src/app.cppm | 27 + tests/cuda-consumer/src/cpu/saxpy.cpp | 11 + tests/cuda-consumer/src/kernels/saxpy.cu | 56 ++ tests/cuda-consumer/src/main.cpp | 12 + tests/spirv-consumer/build.mcpp | 10 + tests/spirv-consumer/mcpp.toml | 36 ++ tests/spirv-consumer/shaders/scale.comp | 15 + tests/spirv-consumer/src/main.cpp | 11 + tools/README.md | 14 + 18 files changed, 1459 insertions(+) create mode 100644 .github/workflows/ci.yml create mode 100644 README.md create mode 100644 mcpp.toml create mode 100644 rules/cuda.cppm create mode 100644 rules/spirv.cppm create mode 100644 src/plugins.cppm create mode 100644 tests/cuda-consumer/build.mcpp create mode 100644 tests/cuda-consumer/include/saxpy/saxpy.h create mode 100644 tests/cuda-consumer/mcpp.toml create mode 100644 tests/cuda-consumer/src/app.cppm create mode 100644 tests/cuda-consumer/src/cpu/saxpy.cpp create mode 100644 tests/cuda-consumer/src/kernels/saxpy.cu create mode 100644 tests/cuda-consumer/src/main.cpp create mode 100644 tests/spirv-consumer/build.mcpp create mode 100644 tests/spirv-consumer/mcpp.toml create mode 100644 tests/spirv-consumer/shaders/scale.comp create mode 100644 tests/spirv-consumer/src/main.cpp create mode 100644 tools/README.md diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml new file mode 100644 index 0000000..88198b4 --- /dev/null +++ b/.github/workflows/ci.yml @@ -0,0 +1,59 @@ +name: ci + +on: + push: + branches: [main] + pull_request: + workflow_dispatch: + +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.3 + +jobs: + consumers: + name: consumers (linux x86_64) + runs-on: ubuntu-24.04 + timeout-minutes: 60 + steps: + - uses: actions/checkout@v4 + + - name: Cache the mcpp sandbox + uses: actions/cache@v4 + with: + path: ~/.mcpp + key: mcpp-sandbox-${{ runner.os }}-${{ runner.arch }}-${{ env.MCPP_VERSION }}-${{ hashFiles('tests/*/mcpp.toml') }} + restore-keys: | + mcpp-sandbox-${{ runner.os }}-${{ runner.arch }}-${{ env.MCPP_VERSION }}- + + - name: Fetch the released mcpp + run: | + curl -L -fsS --retry 3 --retry-all-errors -o mcpp.tar.gz \ + "https://github.com/mcpp-community/mcpp/releases/download/v${MCPP_VERSION}/mcpp-${MCPP_VERSION}-linux-x86_64.tar.gz" + tar -xzf mcpp.tar.gz + MCPP="$PWD/mcpp-${MCPP_VERSION}-linux-x86_64/bin/mcpp" + "$MCPP" --version + "$MCPP" self config --mirror GLOBAL + echo "MCPP=$MCPP" >> "$GITHUB_ENV" + + # The rule's output is a header holding the SPIR-V module; the program + # checks the magic number in its first word, so no Vulkan runtime is + # needed and what is tested is the rule and the engine path feeding it. + - name: rules-spirv through a consumer + working-directory: tests/spirv-consumer + run: | + "$MCPP" build + "$MCPP" run | tee run.log + grep -q '^magic=07230203' run.log + + # Compiles the device unit on a machine with no GPU: the clang route + # produces sm_89 code from the payload toolkit. Running it needs a + # device, so the run is of the CPU variant, which the same seam serves. + - name: rules-cuda through a consumer + working-directory: tests/cuda-consumer + run: | + "$MCPP" build + "$MCPP" build --no-accel + "$MCPP" run --no-accel | tee run.log + grep -q '^12 24 36 48' run.log diff --git a/README.md b/README.md new file mode 100644 index 0000000..fb4b88e --- /dev/null +++ b/README.md @@ -0,0 +1,86 @@ +# mcpp-plugins + +The build plugins the mcpp project maintains, published as one package, +`mcpp:plugins`. A consumer selects the members it needs through features and +imports each one from `build.mcpp` under the module name the member declares. + +```toml +[dependencies.mcpp] +plugins = { version = "0.1.0", features = ["rules-spirv"], host-module = true } +``` + +```cpp +// build.mcpp +import std; +import mcpp; +import mcpp.rules.spirv; + +int main() { + mcpp::rules::spirv::options opt; + opt.includes = { "shaders" }; + return mcpp::rules::spirv::compile(opt) ? 0 : 1; +} +``` + +## Naming + +| family | module name | purpose | +|---|---|---| +| rules | `mcpp.rules.` | how one kind of translation unit is compiled by a compiler mcpp does not drive: the spelling of its flags, the probe of its toolkit, the actions it submits | +| tools | `mcpp.tools.` | a build-time utility independent of any compiler; see `tools/README.md` | +| identity | `mcpp.plugins` | the lib root, compiled before every member; it states the collection's version | + +The `mcpp.` prefix is reserved for this package: mcpp warns when a module under +it is declared by a package outside the `mcpp` namespace. `mcpp.build.*` is the +engine's own module family and is not used here. + +## Members + +| 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`), `[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 | + +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. The index descriptor states the floor; a project +on an older mcpp is refused at resolution rather than at the first shader. + +## How the engine sees this package + +mcpp compiles every module interface unit among a host-module package's +resolved sources as a host module of its own, the lib root first (2026.9.5.3+). +`[features.] sources` is what puts a member into that set, so the module set +a consumer can import is exactly its feature set. A member is compiled alone, +in the same command as the consumer's `build.mcpp`, and may import `std`, +`mcpp` and `mcpp.plugins`. + +## Layout + +``` +mcpp.toml the package: one feature per member +src/plugins.cppm export module mcpp.plugins; +rules/.cppm export module mcpp.rules.; +tools/.cppm export module mcpp.tools.; +tests// one project per member, built by CI with the pinned mcpp +``` + +## Adding a member + +1. One file, `rules/.cppm` or `tools/.cppm`, declaring its module name. + The engine owns the graph — the accelerator axis, the constrained globs, + the action edges, the fingerprint — and the member owns the spelling: which + tool, which flags, what is generated. A member does not read `/usr`; a tool + comes from `mcpp::toolchain_dir()`, `mcpp::xpkg_dir()` or an explicit option, + and a missing one is refused naming the `[xlings.workspace]` entry to add. +2. A feature in `mcpp.toml` adding that one source. +3. A consumer under `tests/` that builds through it and asserts on the + artefact, and a row in the table above with the mcpp floor. +4. A version bump: mcpp identifies an installed package by `(name, version)`, + so a changed payload under an unchanged version is not reinstalled. + +## Releases + +A tag `v` publishes the source archive; the index descriptor +(`mcpp-index/pkgs/m/mcpp.plugins.lua`) names the GitHub archive and its GitCode +mirror with one sha256. diff --git a/mcpp.toml b/mcpp.toml new file mode 100644 index 0000000..739608d --- /dev/null +++ b/mcpp.toml @@ -0,0 +1,28 @@ +[package] +name = "plugins" +namespace = "mcpp" +version = "0.1.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"] +repo = "https://github.com/mcpp-community/mcpp-plugins" + +[language] +standard = "c++23" +modules = true +import_std = true + +# The lib root alone is compiled for a consumer that activates no feature. Each +# member of the collection is a module interface unit that one feature adds to +# the source set; mcpp compiles every interface unit among the resolved +# sources as a host module under the name the unit declares (2026.9.5.3+). +[build] +sources = ["src/plugins.cppm"] + +[features] +default = [] +rules-cuda = { sources = ["rules/cuda.cppm"] } +rules-spirv = { sources = ["rules/spirv.cppm"] } + +[targets.plugins] +kind = "lib" diff --git a/rules/cuda.cppm b/rules/cuda.cppm new file mode 100644 index 0000000..48f0ee1 --- /dev/null +++ b/rules/cuda.cppm @@ -0,0 +1,619 @@ +// 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 tools: where the toolkit's +// pieces live, how an architecture is spelled, which host compilers nvcc +// tolerates, how a driver states its version. None of it is knowledge about +// the build graph. The engine owns the graph, the artifact's identity and the +// architecture set (`[build] accel`); the spelling of the command that +// produces an object, and every probe of the machine, is this file's business. +// `tests/unit/test_core_vendor_probes.cpp` in mcpp holds that line from the +// other side: the engine names no vendor tool. +// +// TWO ROUTES, ONE PRIMARY +// +// `clang -x cuda` is the primary route: the resolved toolchain's own clang +// compiles the device unit, so there is no second host compiler and no +// host-compiler bound to satisfy. nvcc is the alternate, taken when the +// project's toolchain is GCC or when asked for: it drives a host compiler +// (`-ccbin`) and refuses one newer than the bound its `crt/host_config.h` +// states, which this rule reads and reports. +// +// WHAT THE RULE TELLS THE ENGINE +// +// Objects, through `mcpp::action` (role "object"); library directories, +// through `mcpp::link_search`; and three claims about the machine that the +// engine compares or relays before the first compile: +// - the driver's version, read through the driver's own library and stated +// with `mcpp::fact`, together with the floor the runtime needs +// (`mcpp::floor`); an unmet floor refuses the build with both values; +// - whether nvcc can reach its own back-end stages (`--dryrun`), as an +// advisory naming the first stage that does not resolve; +// - whether the embedded PTX can be JIT-compiled by this driver, as an +// advisory, because the SASS for the named architectures still runs. + +module; +#include +#include +#if !defined(_WIN32) +#include +#endif + +export module mcpp.rules.cuda; + +import std; +import mcpp; + +export namespace mcpp::rules::cuda { + +enum class route { automatic, clang, nvcc }; + +struct options { + route which = route::automatic; + // Header search paths for the island. Relative entries resolve against the + // package root; an ABSOLUTE entry is passed through unchanged. + // + // ⭐ THE ABSOLUTE FORM IS FOR A DEPENDENCY'S HEADERS. A device compiler is + // a separate driver and inherits nothing from the C++ side's include + // configuration, so a package whose device code includes a dependency's + // header -- ggml's CUDA backend includes `cublas_v2.h` -- has to name that + // dependency's directory here, and it knows it only as the absolute path + // `mcpp::dep_dir` answered with. + std::vector includes; + std::string out_dir = std::string(mcpp::out_dir()); +}; + +// ─── What the engine said ────────────────────────────────────────────────── + +// The `cuda` chunk of `mcpp::accel()`, in this rule's own reading: the engine +// carries the string and compares it as a shape; what `sm_89` means is ours. +struct target { + std::string version; // "12.9" + std::vector archs; // {"sm_89"} + std::string ptx; // "89" when a portable form is embedded + bool present = false; +}; + +// Split on one character. Written out rather than taken from : GCC 16 +// refuses the ranges split view instantiated inside an exported inline function +// when build.mcpp imports this module (`conflicting deduced return type for +// imported declaration ... view_interface::data()`), and clang does not. +inline std::vector split(std::string_view s, char sep) { + std::vector out; + for (std::size_t i = 0; i <= s.size();) { + auto j = s.find(sep, i); + out.push_back(s.substr(i, j == std::string_view::npos ? s.size() - i : j - i)); + if (j == std::string_view::npos) break; + i = j + 1; + } + return out; +} + +inline std::string_view trim(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; +} + +inline target parse_target(std::string_view accel) { + target t; + for (std::size_t i = 0; i <= accel.size();) { + auto comma = accel.find(',', i); + auto open = accel.find('{', i), close = accel.find('}', i); + if (open != std::string_view::npos && close != std::string_view::npos + && comma != std::string_view::npos && comma > open && comma < close) + comma = accel.find(',', close); + auto chunk = trim(comma == std::string_view::npos ? accel.substr(i) + : accel.substr(i, comma - i)); + i = comma == std::string_view::npos ? accel.size() + 1 : comma + 1; + if (!chunk.starts_with("cuda")) continue; + t.present = true; + auto plus = chunk.find('+'); + t.version = std::string(trim(chunk.substr(4, plus == std::string_view::npos + ? chunk.size() - 4 : plus - 4))); + if (plus != std::string_view::npos) { + auto o = chunk.find('{', plus), c = chunk.find('}', plus); + if (o != std::string_view::npos && c != std::string_view::npos) + for (auto part : split(chunk.substr(o + 1, c - o - 1), ',')) + if (auto a = trim(part); !a.empty()) t.archs.emplace_back(a); + auto tail = chunk.substr(c == std::string_view::npos ? chunk.size() : c + 1); + for (auto key : {"ptx>=", "floor>="}) + if (auto p = tail.find(key); p != std::string_view::npos) + t.ptx = std::string(trim(tail.substr(p + std::string_view(key).size()))); + } + } + return t; +} + +inline std::vector device_sources() { + std::vector out; + for (auto part : split(std::string_view(mcpp::device_sources()), '\n')) + if (auto s = trim(part); !s.empty()) out.emplace_back(s); + return out; +} + +// ─── Locating the toolkit ────────────────────────────────────────────────── + +// The toolkit this project declared under `[xlings.workspace]`, by component. +// The 13.x line splits the compiler across `cuda-nvcc`, `cuda-crt` and +// `libnvvm`; the 12.x line keeps them in `cuda-nvcc`. Either way the project +// names the compiler and this rule finds the pieces. +struct toolkit { + std::string nvcc_root, cudart_root, crt_root, driver_dir; + std::string nvcc() const { return nvcc_root + "/bin/nvcc"; } + std::string host_config() const { + for (auto const* r : { &crt_root, &nvcc_root, &cudart_root }) { + if (r->empty()) continue; + auto p = *r + "/include/crt/host_config.h"; + if (std::filesystem::exists(p)) return p; + } + return {}; + } + std::vector include_dirs() const { + std::vector out; + for (auto const* r : { &cudart_root, &crt_root, &nvcc_root }) + if (!r->empty() && std::filesystem::is_directory(*r + "/include")) + out.push_back(*r + "/include"); + return out; + } + std::vector lib_dirs() const { + std::vector out; + for (auto const* r : { &cudart_root, &nvcc_root }) + for (auto const* sub : { "/lib", "/lib64" }) + if (!r->empty() && std::filesystem::is_directory(*r + sub)) + out.push_back(*r + sub); + return out; + } +}; + +inline std::string xpkg(const char* name) { + if (const char* d = mcpp::xpkg_dir("xim", name); d && *d) return d; + return {}; +} + +inline std::optional find_toolkit() { + toolkit t; + t.nvcc_root = xpkg("cuda-nvcc"); + t.cudart_root = xpkg("cuda-cudart"); + t.crt_root = xpkg("cuda-crt"); + t.driver_dir = xpkg("libcuda-host-link"); + if (t.nvcc_root.empty() || t.cudart_root.empty()) { + std::println(std::cerr, + "mcpp.rules.cuda: the toolkit is not declared.\n" + " Name it under [xlings.workspace] and mcpp provisions it on first use:\n" + " \"xim:cuda-nvcc\" = \"12.9.86\"\n" + " \"xim:cuda-cudart\" = \"12.9.79\"\n" + " (found nvcc: '{}', cudart: '{}')", t.nvcc_root, t.cudart_root); + return std::nullopt; + } + return t; +} + +// ─── Probes: what the machine has, what the toolkit needs ────────────────── + +// The driver's version through the driver's own library, reached through the +// sentinel package rather than /usr/lib. "" when there is no driver here, +// which is a fact about the machine and not a failure of the build. +inline std::string driver_version(const toolkit& t) { +#if defined(_WIN32) + return {}; +#else + if (t.driver_dir.empty()) return {}; + const auto lib = t.driver_dir + "/lib/libcuda.so.1"; + if (!std::filesystem::exists(lib)) return {}; + // What would change the answer is the library the answer was read from. + mcpp::rerun_if_changed(lib.c_str()); + void* h = ::dlopen(lib.c_str(), RTLD_LAZY | RTLD_LOCAL); + if (!h) { + const char* why = ::dlerror(); + mcpp::warning(std::format("could not open the driver library {}: {}", lib, + why ? why : "(no reason given)").c_str()); + return {}; + } + using fn = int (*)(int*); + auto get = reinterpret_cast(::dlsym(h, "cuDriverGetVersion")); + int v = 0; + std::string out; + if (get && get(&v) == 0 && v > 0) out = std::format("{}.{}", v / 1000, (v % 1000) / 10); + ::dlclose(h); + return out; +#endif +} + +inline int major_of(std::string_view v) { + int m = 0; + for (char c : v) { if (!std::isdigit(static_cast(c))) break; m = m * 10 + (c - '0'); } + return m; +} + +inline bool version_at_least(std::string_view have, std::string_view want) { + auto parse = [](std::string_view s) { + std::vector out; int acc = 0; bool digits = false; + for (char c : s) { + if (c == '.') { out.push_back(acc); acc = 0; digits = false; continue; } + if (!std::isdigit(static_cast(c))) break; + acc = acc * 10 + (c - '0'); digits = true; + } + if (digits) out.push_back(acc); + return out; + }; + auto h = parse(have), w = parse(want); + for (std::size_t i = 0; i < std::max(h.size(), w.size()); ++i) { + int a = i < h.size() ? h[i] : 0, b = i < w.size() ? w[i] : 0; + if (a != b) return a > b; + } + return true; +} + +// State the driver relation. The engine compares the floor against the fact +// and refuses with both values; this rule only knows which numbers matter. +// +// The floor is the toolkit's major: a 12.x runtime runs on any 12.x driver +// (minor-version compatibility), and fails at the first allocation on an 11.x +// one. The embedded PTX is a separate, softer question: PTX emitted by toolkit +// 12.9 is JIT-compiled only by a driver at or above 12.9, but the SASS for the +// named architectures still runs, so a driver below the toolkit costs reach on +// newer hardware rather than correctness here -- reported, not enforced. +inline void state_driver_relation(const toolkit& t, const target& tg) { + const auto driver = driver_version(t); + if (!driver.empty()) mcpp::fact("cuda.driver", driver.c_str()); + else mcpp::warning("no driver library reachable through xim:libcuda-host-link; " + "the build proceeds and the artifact will find no device at run time"); + const int major = major_of(tg.version); + if (major > 0) mcpp::floor(std::format("cuda.driver >= {}.0", major).c_str()); + if (!driver.empty() && !tg.ptx.empty() && !version_at_least(driver, tg.version)) + mcpp::warning(std::format( + "the PTX embedded for compute_{} was emitted by toolkit {} and this driver " + "serves {}; hardware newer than {{{}}} will not be able to JIT it. The named " + "architectures run. Build with a toolkit at or below the driver, or add the " + "newer hardware's SASS to [build] accel.", + tg.ptx, tg.version, driver, [&] { + std::string s; for (auto& a : tg.archs) { if (!s.empty()) s += ','; s += a; } + return s; }()).c_str()); +} + +// 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; +} + +// Does the C library this build compiles against declare the C23 functions +// `cospi`, `sinpi` and `rsqrt`? +// +// ⚠️ Measured 2026-09-05 against glibc 2.44. Toolkit 12.9's +// `crt/math_functions.h` declares those same names for the host WITHOUT +// `noexcept`; glibc declares them WITH it, and since C++17 that is part of the +// function type. nvcc's front end stops with six `exception specification is +// incompatible` errors that name a glibc header and a CUDA header and leave +// the reader to work out that neither is at fault alone. The 13.x line does +// not redeclare them and compiles cleanly against the same C library. +// +// Read, not probed. The answer is one substring of one header the sysroot +// already contains; a probe compile would spend a second nvcc invocation to +// learn the same thing, and would report it as a compile failure rather than +// as a pairing that cannot work. +inline bool libc_declares_c23_pi_math(std::string_view sysroot) { + if (sysroot.empty()) return false; + for (auto const* rel : { "/usr/include/bits/mathcalls.h", "/include/bits/mathcalls.h" }) { + std::ifstream in{std::string(sysroot) + rel}; + if (!in) continue; + std::string text{std::istreambuf_iterator(in), std::istreambuf_iterator()}; + return text.find("(cospi,") != std::string::npos + && text.find("(rsqrt,") != std::string::npos; + } + return false; +} + +// ` -dumpversion` → major. The host compiler is the toolchain mcpp +// resolved for this build, so its version is a fact of the build, not a guess. +inline int compiler_major(const std::string& cc) { +#if defined(_WIN32) + (void)cc; return 0; +#else + std::string cmd = cc + " -dumpversion 2>/dev/null"; + if (FILE* p = ::popen(cmd.c_str(), "r")) { + char buf[64] = {}; + std::string s; + if (std::fgets(buf, sizeof buf, p)) s = buf; + ::pclose(p); + return major_of(s); + } + return 0; +#endif +} + +// The first back-end stage nvcc names but cannot resolve, from its own plan. +// nvcc invokes cicc, cudafe++, ptxas and fatbinary by bare name on a PATH it +// states in the plan; a stage that does not resolve there fails the compile +// with `sh: 1: cicc: not found`, naming nothing that helps. +inline std::optional unreachable_stage(const toolkit& t, const std::string& ccbin) { +#if defined(_WIN32) + (void)t; (void)ccbin; return std::nullopt; +#else + const auto probe = std::filesystem::temp_directory_path() / "mcpp-rules-cuda-dryrun.cu"; + { std::ofstream(probe) << "__global__ void k() {}\n"; } + std::string cmd = std::format("{} --dryrun -ccbin {} -c {} -o /dev/null 2>&1", + t.nvcc(), ccbin, probe.string()); + std::string text; + if (FILE* p = ::popen(cmd.c_str(), "r")) { + char buf[4096]; + while (std::fgets(buf, sizeof buf, p)) text += buf; + ::pclose(p); + } + std::filesystem::remove(probe); + std::string path; + std::vector stages; + for (auto l : split(text, '\n')) { + if (!l.starts_with("#$ ")) continue; + l.remove_prefix(3); + if (l.starts_with("PATH=")) { path = std::string(l.substr(5)); continue; } + for (auto const* stage : { "cicc", "cudafe++", "ptxas", "fatbinary", "nvlink" }) { + auto pos = l.find(stage); + if (pos == 0 || (pos != std::string_view::npos && (l[pos - 1] == ' ' || l[pos - 1] == '"'))) + if (std::ranges::find(stages, stage) == stages.end()) stages.emplace_back(stage); + } + } + if (stages.empty()) return std::nullopt; // no plan, no finding + for (auto const& stage : stages) { + bool found = false; + for (auto dir : split(path, ':')) { + std::string d(dir); + if (!d.empty() && std::filesystem::exists(d + "/" + stage)) { found = true; break; } + } + if (!found) return stage; + } + return std::nullopt; +#endif +} + +// ─── Planning ────────────────────────────────────────────────────────────── + +struct edge { + std::string id, description; + std::vector command, inputs, outputs; +}; + +inline route decide(route asked) { + if (asked != route::automatic) return asked; + return std::string_view(mcpp::compiler()) == "clang" ? route::clang : route::nvcc; +} + +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, "mcpp.rules.cuda: no mcpp build context -- this runs from build.mcpp"); + return out; + } + const auto tg = parse_target(mcpp::accel()); + if (!tg.present || tg.archs.empty()) { + // C19: a device build that names no device is refused HERE, not at + // run time as `no kernel image is available for execution`. + std::println(std::cerr, + "mcpp.rules.cuda: [build] accel names no CUDA architecture (accel = \"{}\").\n" + " Write e.g. accel = \"cuda12.9+{{sm_89}} ptx>=89\" -- the set a build compiles\n" + " for is a decision, and the machine's own hardware is a poor default for it.", + mcpp::accel()); + return out; + } + auto tk = find_toolkit(); + if (!tk) return out; + state_driver_relation(*tk, tg); + + const route r = decide(opt.which); + const std::string tcdir = mcpp::toolchain_dir(); + std::string driver_cc; // the compiler that runs the device unit + std::vector front; // the command up to the input file + if (r == route::clang) { + driver_cc = tcdir + "/bin/clang++"; + if (!std::filesystem::exists(driver_cc)) { + std::println(std::cerr, "mcpp.rules.cuda: the clang route needs the toolchain's clang++ at {}", driver_cc); + return out; + } + front = { driver_cc, "-x", "cuda", "-std=c++17", "-O2", "-fPIC", + "--cuda-path=" + tk->nvcc_root, "-Wno-unknown-cuda-version", + // ⚠️ NVIDIA'S HEADER REFUSES libc++, AND THE REFUSAL IS + // ABOUT nvcc RATHER THAN ABOUT THIS COMPILER. + // + // crt/host_defines.h:67: error: "libc++ is not supported + // on x86 system" + // + // The guard is `#if defined(__CUDACC__) && … && + // defined(_LIBCPP_VERSION)`, and clang defines `__CUDACC__` + // when it compiles CUDA itself — so a device unit that + // includes stops here on any LLVM + // toolchain, which is the toolchain this route exists for. + // Measured on ggml's CUDA backend; the CUDA example's own + // kernel never showed it because a bare kernel includes no + // toolkit header at all. + // + // The escape hatch is upstream's own, and it is passed only + // on this route: nvcc's host pass really does break against + // libc++, and nothing here weakens that. + "-D_ALLOW_UNSUPPORTED_LIBCPP" }; + for (auto const& inc : tk->include_dirs()) front.push_back("-I" + inc); + for (auto const& a : tg.archs) front.push_back("--cuda-gpu-arch=" + a); + // clang checks ptxas and fatbinary itself; say so before it does. + for (auto const* tool : { "ptxas", "fatbinary" }) + if (!std::filesystem::exists(tk->nvcc_root + "/bin/" + tool)) + mcpp::warning(std::format("the toolkit payload has no {}; clang invokes it " + "after generating PTX", tool).c_str()); + std::println("mcpp.rules.cuda: clang route -- {} (toolkit {})", driver_cc, tk->nvcc_root); + } else { + // nvcc drives the toolchain's own compiler, and refuses one newer than + // the bound its header states. Read the bound; if exceeded, pass the + // escape hatch and say so -- an unexplained flag is worse than a note. + const bool clangHost = std::string_view(mcpp::compiler()) == "clang"; + if (clangHost) { + // Measured: nvcc's own crt/host_defines.h stops the compile with + // `libc++ is not supported on x86 system`, and libc++ is what an + // LLVM toolchain's clang uses. The pairing that works is nvcc with + // a GCC toolchain; with an LLVM toolchain the clang route is the + // one to take, and it is the default. + std::println(std::cerr, + "mcpp.rules.cuda: the nvcc route needs a GCC host compiler; this project's " + "toolchain is LLVM, whose clang uses libc++ and nvcc refuses it. Use the clang " + "route (the default for an LLVM toolchain) or set [toolchain] to a gcc payload."); + return out; + } + // The other pairing this route cannot have: an old toolkit and a C + // library new enough to have the C23 `pi` functions. Stated before the + // compile, because the compile's own report names two headers and no + // decision. + if (major_of(tg.version) < 13 + && libc_declares_c23_pi_math(mcpp::toolchain_sysroot())) { + std::println(std::cerr, + "mcpp.rules.cuda: toolkit {} redeclares the C23 functions cospi, sinpi and " + "rsqrt for the host without `noexcept`, and the C library this build compiles " + "against declares them with it; nvcc's front end refuses the pair.\n" + " Name a 13.x toolkit, whose headers leave them to the C library:\n" + " [xlings.workspace]\n" + " \"xim:cuda-nvcc\" = \"13.3.33\"\n" + " \"xim:cuda-crt\" = \"13.3.33\"\n" + " \"xim:cuda-cudart\" = \"13.3.29\"\n" + " or take the clang route, which does not include that header at all.", + tg.version); + return out; + } + // The host compiler nvcc drives, chosen within the bound the toolkit + // states. Measured: gcc 16 under nvcc 12.9 (bound gcc <= 14) fails inside + // nvcc's front end on GCC 16's own even with + // -allow-unsupported-compiler -- the escape hatch admits a compiler one + // step past the bound, not a standard library two majors newer. So the + // rule does not guess: the toolchain's g++ when it is within the bound, + // otherwise a gcc payload the project declared for this purpose, and + // otherwise a refusal that says which declaration to add. + const auto b = read_bounds(tk->host_config()); + const std::string tcGcc = tcdir + "/bin/g++"; + const int tcMajor = compiler_major(tcGcc); + if (b.gcc == 0 || tcMajor <= b.gcc) { + driver_cc = tcGcc; + } else if (auto payload = xpkg("gcc"); !payload.empty() + && compiler_major(payload + "/bin/g++") <= b.gcc) { + driver_cc = payload + "/bin/g++"; + mcpp::warning(std::format( + "nvcc {} states gcc <= {} in {}; the toolchain's gcc {} exceeds it, so the " + "device unit is compiled with the declared xim:gcc payload ({}). The clang " + "route has no such bound.", tg.version, b.gcc, tk->host_config(), tcMajor, + driver_cc).c_str()); + } else { + std::println(std::cerr, + "mcpp.rules.cuda: nvcc {} accepts gcc <= {} ({}), and this project's " + "toolchain is gcc {}.\n" + " Declare a gcc payload within the bound and the rule drives that one:\n" + " [xlings.workspace]\n" + " \"xim:gcc\" = \"13.3.0\"\n" + " or take the clang route with [toolchain] default = \"llvm@22.1.8\".", + tg.version, b.gcc, tk->host_config(), tcMajor); + return out; + } + front = { tk->nvcc(), "-ccbin", driver_cc, "-std=c++17", "-O2", + "--compiler-options", "-fPIC" }; + // The host compiler nvcc drives is not one mcpp resolved, so nothing + // has told it where the C library or the assembler are. Measured: with + // neither of these, NVIDIA's own crt/host_config.h stops at + // `features.h: No such file or directory`. Both are the flags mcpp + // passes to its own compiler for this target. + if (const char* sr = mcpp::toolchain_sysroot(); sr && *sr) { + front.push_back("--compiler-options"); + front.push_back(std::string("--sysroot=") + sr); + } + if (const char* bu = mcpp::toolchain_binutils_dir(); bu && *bu) { + front.push_back("--compiler-options"); + front.push_back(std::string("-B") + bu); + } + for (auto const& inc : tk->include_dirs()) front.push_back("-I" + inc); + for (auto const& a : tg.archs) { + std::string digits; + for (char c : a) if (std::isdigit(static_cast(c))) digits += c; + front.push_back("-gencode"); + front.push_back(std::format("arch=compute_{},code={}", digits, a)); + } + if (!tg.ptx.empty()) { + front.push_back("-gencode"); + front.push_back(std::format("arch=compute_{0},code=compute_{0}", tg.ptx)); + } + if (auto missing = unreachable_stage(*tk, driver_cc)) + mcpp::warning(std::format( + "nvcc cannot reach its own back-end: it invokes '{}' by name and that name " + "does not resolve on the search path it states. On the 13.x line install " + "xim:libnvvm beside xim:cuda-nvcc.", *missing).c_str()); + std::println("mcpp.rules.cuda: nvcc route -- {} with -ccbin {}", tk->nvcc(), driver_cc); + } + + // The link line gets its directories from here, not from the manifest: the + // rule resolved the payload, so the rule names where its libraries are. + for (auto const& d : tk->lib_dirs()) mcpp::link_search(d.c_str()); + + 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 = (r == route::clang ? "clang -x cuda " : "nvcc ") + src; + e.command = front; + for (auto const& inc : opt.includes) + e.command.push_back("-I" + (std::filesystem::path(inc).is_absolute() + ? inc : root + "/" + inc)); + e.command.insert(e.command.end(), { "-c", root + "/" + src, "-o", obj }); + e.inputs = { root + "/" + src }; + e.outputs = { obj }; + out.push_back(std::move(e)); + } + return out; +} + +inline bool submit(std::span edges) { + for (auto const& e : edges) { + mcpp::action a; + a.id = e.id.c_str(); + a.role = "object"; // the linkable artifact itself + 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; +} + +// Everything from the manifest: the architectures from `[build] accel`, the +// sources from the constrained glob in `[build] sources`. A build that asks +// for no accelerator has no device sources and nothing to do here -- that is +// the CPU-only variant, and the seam's fallback carries it. +inline bool compile(options opt = {}) { + if (!*mcpp::accel()) return true; + const auto sources = device_sources(); + if (sources.empty()) { + mcpp::warning("[build] accel names a device but no constrained glob matched a device " + "source; nothing was compiled for it"); + return true; + } + auto edges = plan(sources, std::move(opt)); + if (edges.empty()) return false; + return submit(edges); +} + +} // namespace mcpp::rules::cuda diff --git a/rules/spirv.cppm b/rules/spirv.cppm new file mode 100644 index 0000000..3e8538f --- /dev/null +++ b/rules/spirv.cppm @@ -0,0 +1,356 @@ +// mcpp.rules.spirv — how a GLSL translation unit becomes SPIR-V, stated once. +// +// THE MODULE NAME IS `mcpp.rules.` AND THE PACKAGE NAMESPACE IS `mcpp`. +// Both halves are the rule-package specification, not a preference. The module +// name is declared by this source rather than derived from the package name +// (spec I1), and `mcpp.*` is reserved for rules the mcpp project maintains +// (I8) -- which is enforced as a warning keyed on the package NAMESPACE, so a +// rule carrying this module name under any other namespace is told that it +// claims an origin it does not have. +// +// The division of labour is the one `mcpp.rules.cuda` established, and it is the +// point of both packages: the ENGINE owns the graph — the accelerator axis, +// the constrained source globs that route `shaders/*.comp` here instead of to +// the C++ compiler, the action edges and their ordering, the fingerprint — and +// does not know the word "vulkan" or the word "glslang". The RULE owns the +// spelling: which compiler, which flags, what the generated symbol is called. +// +// WHAT COMES OUT IS A HEADER, NOT AN OBJECT. +// +// A SPIR-V module is data the program hands to `vkCreateShaderModule`, not +// code the linker places. Two shapes are possible: a `.spv` file beside the +// binary, which makes the program's correctness depend on its working +// directory, or a C array compiled into it. This rule emits the second, so the +// artifact carries its shaders and a `mcpp pack` of it has nothing further to +// collect. +// +// `role = "source"` is what makes that work. It is the one role the engine +// orders BEFORE compilation (`action_precedes_compilation`), which is exactly +// what a generated header needs and exactly what an `artifact` role would not +// give: an artifact output is ordered against the LINK, and the header has to +// exist before the first translation unit that includes it is compiled. +// +// GLSLANG, AND ONLY GLSLANG, DELIBERATELY. +// +// `glslc` (shaderc) is the other reference compiler and is not supported here. +// Not for a reason of principle — this rule would take it — but because +// nothing in this ecosystem publishes it, and a route with no payload behind +// it is a claim rather than a feature. `xim:glslang` exists and is what the +// example installs. If glslc is ever packaged, `-mfmt=c` emits a bare +// initialiser list where glslang's `-x --vn` emits a complete declaration, so +// the two produce different headers and the rule would have to say which. +module; +#include +#include + +export module mcpp.rules.spirv; + +import std; +import mcpp; + +export namespace mcpp::rules::spirv { + +struct options { + // The Vulkan environment the SPIR-V targets. Left empty, it is taken from + // the accelerator axis — `accel = "vulkan1.2"` in the manifest — which is + // the same route `mcpp.rules.cuda` takes for `sm_89`, and the reason neither + // rule needs a second place to state what the build is for. + std::string target_env; + // `-I` for GLSL `#include`, `-D` for its preprocessor. Relative entries + // resolve against the package root; an absolute entry is passed through. + std::vector includes; + std::vector defines; + // glslang's optimiser (`-Os`), which is spirv-opt linked into it. + bool optimize = true; + // An explicit compiler path wins over discovery. Set it when a project + // pins a glslang other than the one the workspace installed. + std::string compiler; + std::string out_dir = std::string(mcpp::out_dir()); +}; + +// Where the generated headers are written, and what the build program passes +// to `mcpp::include_dir` so `#include "scale_comp.h"` resolves. +inline std::string include_dir(const options& opt) { + return (std::filesystem::path(opt.out_dir) / "spirv").string(); +} + +// ─── What the engine said ────────────────────────────────────────────────── + +// The `vulkan` chunk of `mcpp::accel()`. Unlike CUDA's, it carries no +// architecture set: SPIR-V is the portable form, and which GPU executes it is +// decided when the driver compiles it, not here. A rule that demanded an +// architecture would be inventing a requirement its device API does not have. +struct target { + std::string version; // "1.2", from `vulkan1.2` + bool present = false; +}; + +inline std::string_view trim(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; +} + +inline target parse_target(std::string_view accel) { + target t; + for (std::size_t i = 0; i <= accel.size();) { + auto comma = accel.find(',', i); + auto chunk = trim(comma == std::string_view::npos ? accel.substr(i) + : accel.substr(i, comma - i)); + i = comma == std::string_view::npos ? accel.size() + 1 : comma + 1; + if (!chunk.starts_with("vulkan")) continue; + t.present = true; + auto rest = chunk.substr(std::string_view("vulkan").size()); + // A bare `vulkan` is legitimate and means "whatever the loader offers"; + // the target environment then falls back to the default below. + auto plus = rest.find('+'); + t.version = std::string(trim(plus == std::string_view::npos ? rest + : rest.substr(0, plus))); + } + return t; +} + +// ─── The compiler ────────────────────────────────────────────────────────── + +inline bool is_file(const std::string& p) { + std::error_code ec; + return !p.empty() && std::filesystem::is_regular_file(p, ec); +} + +// Discovery, in the order a project can predict: what it named, what the +// environment named, the payload the workspace installed, then the PATH. The +// PATH comes last on purpose — a host glslang is a fine fallback and a poor +// default, because it makes the SPIR-V depend on a machine rather than on a +// declaration. +inline std::string find_compiler(const options& opt) { + if (!opt.compiler.empty()) return opt.compiler; + if (const char* e = std::getenv("MCPP_GLSLANG"); e && *e) return e; + + if (const char* dir = mcpp::xpkg_dir("glslang"); dir && *dir) { + for (const char* exe : {"glslangValidator", "glslang"}) { + auto p = (std::filesystem::path(dir) / "bin" / exe).string(); + if (is_file(p)) return p; + } + } + for (const char* exe : {"glslangValidator", "glslang"}) { + if (const char* path = std::getenv("PATH"); path && *path) { + std::string_view sv(path); + for (std::size_t i = 0; i <= sv.size();) { + auto sep = sv.find(':', i); + auto dir = sv.substr(i, sep == std::string_view::npos ? sv.size() - i : sep - i); + i = sep == std::string_view::npos ? sv.size() + 1 : sep + 1; + if (dir.empty()) continue; + auto p = (std::filesystem::path(dir) / exe).string(); + if (is_file(p)) return p; + } + } + } + return {}; +} + +// `Glslang Version: 11:15.1.0` — the first field is the SPIR-V generator +// magic, the second is the release. The release is what a floor compares, and +// stating it as a fact is what makes a build log answer "which compiler +// produced this SPIR-V" without anyone having to reproduce the build. +inline std::string run_and_capture(const std::string& cmd) { + FILE* p = ::popen(cmd.c_str(), "r"); + if (!p) return {}; + std::string text; + char buf[512]; + while (std::fgets(buf, sizeof buf, p)) text += buf; + ::pclose(p); + return text; +} + +// ⚠️ THE OPTIMISER IS OPTIONAL AND ITS ABSENCE IS NOT A BUILD ERROR. +// +// glslang links spirv-opt only when built with `ENABLE_OPT`, and the payload +// this ecosystem publishes today is not: +// +// glslangValidator: Error: -Os not available; optimizer not linked +// +// Passing `-Os` to such a binary fails the compile. A shader that is +// unoptimised is still a correct shader, so the flag is dropped and the build +// says so once — refusing would make an optional pass a requirement, and +// passing it silently would let the manifest claim an optimisation that did +// not happen. +inline bool has_optimizer(const std::string& exe) { + const auto out = run_and_capture("\"" + exe + "\" -Os --version 2>&1"); + return out.find("optimizer not linked") == std::string::npos; +} + +inline std::string compiler_version(const std::string& exe) { + const std::string text = run_and_capture("\"" + exe + "\" --version 2>/dev/null"); + for (std::size_t i = 0; i <= text.size();) { + auto nl = text.find('\n', i); + std::string_view line(text.data() + i, + (nl == std::string::npos ? text.size() : nl) - i); + i = nl == std::string::npos ? text.size() + 1 : nl + 1; + auto at = line.find("Glslang Version:"); + if (at == std::string_view::npos) continue; + auto rest = trim(line.substr(at + std::string_view("Glslang Version:").size())); + auto colon = rest.find(':'); + if (colon != std::string_view::npos) rest = rest.substr(colon + 1); + while (!rest.empty() && (rest.back() == '\n' || rest.back() == '\r')) + rest.remove_suffix(1); + return std::string(trim(rest)); + } + return {}; +} + +// ─── Shaders ─────────────────────────────────────────────────────────────── + +// glslang infers the stage from the extension, and so does this table — but +// the table is consulted rather than trusted: an extension that is not a stage +// is refused by name instead of being handed to a compiler that will refuse it +// with a less specific message. +inline std::string_view stage_of(std::string_view ext) { + if (ext == ".comp") return "comp"; + if (ext == ".vert") return "vert"; + if (ext == ".frag") return "frag"; + if (ext == ".geom") return "geom"; + if (ext == ".tesc") return "tesc"; + if (ext == ".tese") return "tese"; + if (ext == ".mesh") return "mesh"; + if (ext == ".task") return "task"; + if (ext == ".rgen") return "rgen"; + if (ext == ".rint") return "rint"; + if (ext == ".rahit") return "rahit"; + if (ext == ".rchit") return "rchit"; + if (ext == ".rmiss") return "rmiss"; + if (ext == ".rcall") return "rcall"; + return {}; +} + +// `shaders/scale.comp` -> `scale_comp_spv`, and the header that declares it is +// `scale_comp.h`. Derived rather than configurable: a name a project chooses +// per shader is a name the project has to keep in agreement with its own +// `#include`, and this rule already decides the file name. +inline std::string symbol_of(std::string_view stem, std::string_view stage) { + std::string s; + for (char c : stem) + s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; + s += '_'; + s += stage; + s += "_spv"; + return s; +} + +// ⚠️ NEWLINE-SEPARATED, not `;`. A path may contain a semicolon and cannot +// contain a newline, which is why the engine chose it — and why a splitter +// that guesses wrong still works for exactly one shader and silently produces +// one impossible path for two. +inline std::vector device_shaders() { + std::vector out; + std::string_view all(mcpp::device_sources()); + for (std::size_t i = 0; i <= all.size();) { + auto sep = all.find('\n', i); + auto one = trim(all.substr(i, sep == std::string_view::npos ? all.size() - i : sep - i)); + i = sep == std::string_view::npos ? all.size() + 1 : sep + 1; + if (!one.empty()) out.emplace_back(one); + } + return out; +} + +// ─── The rule ────────────────────────────────────────────────────────────── + +inline bool compile(std::span shaders, options opt = {}) { + if (shaders.empty()) return true; + + const auto exe = find_compiler(opt); + if (exe.empty()) { + std::println(stderr, + "mcpp.rules.spirv: no glslang found. Install one into the workspace\n" + " [xlings.workspace]\n" + " \"xim:glslang\" = \"15.1.0\"\n" + "or name it: MCPP_GLSLANG=/path/to/glslangValidator, or set " + "options::compiler."); + return false; + } + if (const auto v = compiler_version(exe); !v.empty()) + mcpp::fact("glslang", v.c_str()); + + bool optimize = opt.optimize; + if (optimize && !has_optimizer(exe)) { + mcpp::warning("mcpp.rules.spirv: this glslang was built without spirv-opt " + "(-Os not available; optimizer not linked); shaders are " + "compiled unoptimised"); + optimize = false; + } + + auto env = opt.target_env; + if (env.empty()) { + const auto t = parse_target(mcpp::accel()); + env = t.version.empty() ? "vulkan1.0" : "vulkan" + t.version; + } + + const std::string root = mcpp::manifest_dir(); + const auto gen = include_dir(opt); + std::error_code ec; + std::filesystem::create_directories(gen, ec); + + for (auto const& src : shaders) { + const std::filesystem::path p(src); + const auto stage = stage_of(p.extension().string()); + if (stage.empty()) { + std::println(stderr, + "mcpp.rules.spirv: {} has no shader stage. glslang derives the stage from " + "the extension; rename it to one of .comp .vert .frag .geom .tesc " + ".tese .mesh .task .rgen .rint .rahit .rchit .rmiss .rcall", src); + return false; + } + const auto sym = symbol_of(p.stem().string(), stage); + const auto header = (std::filesystem::path(gen) + / (p.stem().string() + "_" + std::string(stage) + ".h")).string(); + const auto input = std::filesystem::path(src).is_absolute() + ? src : root + "/" + src; + + // ⚠️ `id` and `description` are raw pointers the action reads at + // `submit()`; `arg`/`input`/`output` copy, these two do not. Held in + // named strings for the life of the statement that submits. + const std::string id = "spirv:" + src; + const std::string desc = "glslang " + src; + + mcpp::action a; + a.id = id.c_str(); + a.role = "source"; // a header: ordered before compilation + a.description = desc.c_str(); + a.arg(exe.c_str()); + a.arg("-V"); + a.arg("--target-env"); a.arg(env.c_str()); + a.arg("-S"); a.arg(std::string(stage).c_str()); + if (optimize) a.arg("-Os"); + for (auto const& d : opt.defines) a.arg(("-D" + d).c_str()); + for (auto const& i : opt.includes) + a.arg(("-I" + (std::filesystem::path(i).is_absolute() ? i : root + "/" + i)).c_str()); + // `-x --vn` is what makes the output a C declaration rather than a + // binary: a `const uint32_t []` the program includes. + a.arg("-x"); + a.arg("--vn"); a.arg(sym.c_str()); + a.arg("-o"); a.arg(header.c_str()); + a.arg(input.c_str()); + a.input(input.c_str()); + a.output(header.c_str()); + a.submit(); + } + + mcpp::include_dir(gen.c_str()); + return true; +} + +// The whole manifest's worth: the shaders the constrained glob routed here. +// A build that names no accelerator has none, and the seam's CPU side carries +// the program — the same shape `mcpp::rules::cuda::compile()` has, for the same +// reason. +inline bool compile(options opt = {}) { + if (!*mcpp::accel()) return true; + const auto shaders = device_shaders(); + if (shaders.empty()) { + mcpp::warning("[build] accel names vulkan but no constrained glob matched a " + "shader; nothing was compiled for it"); + return true; + } + return compile(std::span(shaders), std::move(opt)); +} + +} // namespace mcpp::rules::spirv diff --git a/src/plugins.cppm b/src/plugins.cppm new file mode 100644 index 0000000..6cc8b02 --- /dev/null +++ b/src/plugins.cppm @@ -0,0 +1,16 @@ +// mcpp.plugins: the identity unit of the collection. +// +// Every member of this package is a module interface unit under rules/ or +// tools/, compiled as a host module of its own when the consumer's feature +// request names it. This unit is the lib root. It is compiled before every +// member, so a member may import it, and it states the one fact a member may +// want to report about itself: the version of the collection it belongs to. +export module mcpp.plugins; + +import std; + +export namespace mcpp::plugins { + +inline constexpr std::string_view version = "0.1.0"; + +} // namespace mcpp::plugins diff --git a/tests/cuda-consumer/build.mcpp b/tests/cuda-consumer/build.mcpp new file mode 100644 index 0000000..1d23a33 --- /dev/null +++ b/tests/cuda-consumer/build.mcpp @@ -0,0 +1,9 @@ +import std; +import mcpp; +import mcpp.rules.cuda; + +int main() { + mcpp::rules::cuda::options opt; + opt.includes = { "include" }; + return mcpp::rules::cuda::compile(opt) ? 0 : 1; +} diff --git a/tests/cuda-consumer/include/saxpy/saxpy.h b/tests/cuda-consumer/include/saxpy/saxpy.h new file mode 100644 index 0000000..8d82d97 --- /dev/null +++ b/tests/cuda-consumer/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/tests/cuda-consumer/mcpp.toml b/tests/cuda-consumer/mcpp.toml new file mode 100644 index 0000000..03d9f46 --- /dev/null +++ b/tests/cuda-consumer/mcpp.toml @@ -0,0 +1,73 @@ +[package] +name = "cuda-consumer" +namespace = "example" +version = "0.1.0" +description = "A CUDA kernel behind a seam module, with a CPU fallback" +accelerators = ["cuda"] + +[language] +standard = "c++23" +modules = true +import_std = true + +# The primary route is clang: the toolchain's own clang++ compiles the device +# unit (`-x cuda`), so there is no second host compiler and no host-compiler +# bound to satisfy. With a GCC toolchain the rule package takes the nvcc route +# instead, driving that GCC and reading the bound nvcc states for it. +[toolchain] +default = "llvm@22.1.8" + +[dependencies.mcpp] +plugins = { path = "../..", features = ["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. It is the one CUDA component +# that cannot be an ordinary payload -- the licence forbids redistributing it +# and it is in ABI lockstep with the kernel module. +[dependencies.compat] +cuda-runtime = "2026.09.05" + +# The toolkit this project builds with, named rather than discovered. These +# are PAYLOADS, so the version is the project's choice and not the machine's. +# The 12.9 line is named on purpose: a runtime must not be newer than the +# driver it will meet, and 12.x reaches every driver from r525 onward. The +# rule package states the driver relation and mcpp compares it before the +# first compile. +[xlings.workspace] +"xim:cuda-nvcc" = "12.9.86" +"xim:cuda-cudart" = "12.9.79" +"xim:libcuda-host-link" = { linux = "0.0.1" } + +[build] +# What this build compiles device code FOR. Written once, here: the rule +# package derives its own flags (`--cuda-gpu-arch`, `-gencode`) from it. +accel = "cuda12.9+{sm_89} ptx>=89" +# The device sources carry the accel they are for. Under `--no-accel` the glob +# is left out and the CPU fallback below takes its place; under an accel that +# does not cover it the build is refused naming both. +sources = [ + "src/*.cppm", + "src/*.cpp", + { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }, +] +include_dirs = ["include"] + +# The CUDA runtime is linked STATICALLY, and only when a device build asks for +# it. Linking the redistributable half in leaves exactly one host dependency, +# libcuda.so.1, which the driver owns and the sentinel package reaches. +# ⭐ NO ABSOLUTE PATHS: the rule package puts the payload's library directory on +# the link line from `mcpp::xpkg_dir`, so this manifest names libraries only. +[target.'cfg(accelerator = "cuda")'.build] +ldflags = ["-lcudart_static", "-lrt", "-lpthread", "-ldl"] + +# The CPU-only variant: the same seam, a host implementation behind it. This +# is what `mcpp build --no-accel` produces, and what a machine with no device +# runs. +[target.'cfg(not(accelerator = "cuda"))'.build] +sources = ["src/cpu/*.cpp"] + +[targets.cuda-consumer] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/cuda-consumer/src/app.cppm b/tests/cuda-consumer/src/app.cppm new file mode 100644 index 0000000..ab6d4dd --- /dev/null +++ b/tests/cuda-consumer/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/tests/cuda-consumer/src/cpu/saxpy.cpp b/tests/cuda-consumer/src/cpu/saxpy.cpp new file mode 100644 index 0000000..4de4ad6 --- /dev/null +++ b/tests/cuda-consumer/src/cpu/saxpy.cpp @@ -0,0 +1,11 @@ +// The CPU implementation behind the same seam. Compiled only when the build +// asks for no accelerator (`mcpp build --no-accel`), through the +// `cfg(not(accelerator = "cuda"))` section of the manifest; the device island +// and this file define the same symbol and are never in one link. +#include "saxpy/saxpy.h" + +extern "C" int saxpy_device(float a, const float* x, const float* y, + float* out, unsigned n) { + for (unsigned i = 0; i < n; ++i) out[i] = a * x[i] + y[i]; + return 0; +} diff --git a/tests/cuda-consumer/src/kernels/saxpy.cu b/tests/cuda-consumer/src/kernels/saxpy.cu new file mode 100644 index 0000000..d5d858d --- /dev/null +++ b/tests/cuda-consumer/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/tests/cuda-consumer/src/main.cpp b/tests/cuda-consumer/src/main.cpp new file mode 100644 index 0000000..6efc193 --- /dev/null +++ b/tests/cuda-consumer/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/tests/spirv-consumer/build.mcpp b/tests/spirv-consumer/build.mcpp new file mode 100644 index 0000000..9f286ae --- /dev/null +++ b/tests/spirv-consumer/build.mcpp @@ -0,0 +1,10 @@ +import std; +import mcpp; +import mcpp.rules.spirv; + +int main() { + mcpp::rerun_if_changed_glob("shaders/**/*.comp"); + mcpp::rules::spirv::options opt; + opt.includes = { "shaders" }; + return mcpp::rules::spirv::compile(opt) ? 0 : 1; +} diff --git a/tests/spirv-consumer/mcpp.toml b/tests/spirv-consumer/mcpp.toml new file mode 100644 index 0000000..ee0d317 --- /dev/null +++ b/tests/spirv-consumer/mcpp.toml @@ -0,0 +1,36 @@ +# The consumer that tests mcpp.rules.spirv. +# +# No Vulkan runtime is involved: the rule's output is a header holding the +# SPIR-V module as a `const uint32_t` array, and the program checks the one +# property that identifies a SPIR-V module without executing it, the magic +# number 0x07230203 in its first word. What is tested is therefore the rule +# and the engine path that feeds it -- the constrained glob, the device +# source list, the `role = "source"` action ordered before compilation -- +# and not a driver. +[package] +name = "spirv-consumer" +namespace = "example" +version = "0.1.0" +description = "A GLSL compute shader compiled to a SPIR-V header through mcpp.rules.spirv" + +[language] +standard = "c++23" +modules = true +import_std = true + +[dependencies.mcpp] +plugins = { path = "../..", features = ["rules-spirv"], host-module = true } + +[xlings.workspace] +"xim:glslang" = "15.1.0" + +[build] +accel = "vulkan1.2" +sources = [ + "src/*.cpp", + { glob = "shaders/*.comp", accel = "vulkan1.2" }, +] + +[targets.spirv-consumer] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/spirv-consumer/shaders/scale.comp b/tests/spirv-consumer/shaders/scale.comp new file mode 100644 index 0000000..360eb26 --- /dev/null +++ b/tests/spirv-consumer/shaders/scale.comp @@ -0,0 +1,15 @@ +#version 450 + +// The device side of the same computation `examples/09-cuda-kernel` runs on +// CUDA: out = a*x + y. One storage buffer holds all three vectors so the host +// side needs one allocation and one descriptor. +layout(local_size_x = 64) in; + +layout(std430, binding = 0) buffer Data { float v[]; }; +layout(push_constant) uniform Push { float a; uint n; } push; + +void main() { + const uint i = gl_GlobalInvocationID.x; + if (i >= push.n) return; + v[2u * push.n + i] = push.a * v[i] + v[push.n + i]; +} diff --git a/tests/spirv-consumer/src/main.cpp b/tests/spirv-consumer/src/main.cpp new file mode 100644 index 0000000..55675ba --- /dev/null +++ b/tests/spirv-consumer/src/main.cpp @@ -0,0 +1,11 @@ +// The generated header declares `const uint32_t scale_comp_spv[]`; the name +// is the shader's stem, its stage, and `_spv`, as mcpp.rules.spirv documents. +#include +#include +#include "scale_comp.h" + +int main() { + const std::uint32_t magic = scale_comp_spv[0]; + std::printf("magic=%08x words=%zu\n", magic, sizeof scale_comp_spv / sizeof scale_comp_spv[0]); + return magic == 0x07230203u ? 0 : 1; +} diff --git a/tools/README.md b/tools/README.md new file mode 100644 index 0000000..391030d --- /dev/null +++ b/tools/README.md @@ -0,0 +1,14 @@ +# mcpp.tools.* + +Build-time utilities that are not rules. A rule (`mcpp.rules.`) states how +one kind of translation unit is compiled by a compiler mcpp does not drive; a +tool states something a build program needs that is independent of any +compiler: generating a header from a data file, computing a value that several +rules share, checking an invariant of the source tree. + +The directory is empty in 0.1.0. A member is added when a second consumer +needs it; a utility written for one consumer belongs in that consumer's +`build.mcpp`. Each member is one file, `tools/.cppm`, declaring +`export module mcpp.tools.;`, added to the source set by the feature +`tools-` in `mcpp.toml`, and importing nothing but `std`, `mcpp` and +`mcpp.plugins`. From 665fa48879413835c2c2b619294129cdd6c00014 Mon Sep 17 00:00:00 2001 From: speak-agent Date: Sat, 5 Sep 2026 20:30:34 +0800 Subject: [PATCH 2/3] fix(rules-cuda): the clang route names the cuRAND payload it needs clang's own CUDA wrapper includes curand_mtgp32_kernel.h unconditionally. On a developer machine the header was found in the host's /usr/include and the dependency went unnoticed; the CI runner, with no host CUDA, refused it. The rule now adds xim:libcurand's include directory to the device compile and refuses the clang route without it, naming the [xlings.workspace] entry to add. The consumer fixture declares it. --- README.md | 2 +- rules/cuda.cppm | 23 ++++++++++++++++++++++- tests/cuda-consumer/mcpp.toml | 4 ++++ 3 files changed, 27 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index fb4b88e..0402ed5 100644 --- a/README.md +++ b/README.md @@ -38,7 +38,7 @@ 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`), `[build] accel = "cuda…"`, a constrained glob for `*.cu`; the clang route with an LLVM toolchain, the nvcc route with a GCC one | +| `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 | The floor is the mcpp release whose engine carries what the member relies on: diff --git a/rules/cuda.cppm b/rules/cuda.cppm index 48f0ee1..534c678 100644 --- a/rules/cuda.cppm +++ b/rules/cuda.cppm @@ -141,6 +141,13 @@ inline std::vector device_sources() { // names the compiler and this rule finds the pieces. struct toolkit { std::string nvcc_root, cudart_root, crt_root, driver_dir; + // cuRAND's headers. Not a runtime dependency of a kernel that never calls + // cuRAND: clang's own CUDA wrapper (`__clang_cuda_runtime_wrapper.h`) + // includes `curand_mtgp32_kernel.h` unconditionally, so the clang route + // cannot compile any device unit without them. Measured 2026-09-05: on a + // developer machine the header was found in the HOST's /usr/include and + // the leak went unnoticed until a runner with no host CUDA refused it. + std::string curand_root; std::string nvcc() const { return nvcc_root + "/bin/nvcc"; } std::string host_config() const { for (auto const* r : { &crt_root, &nvcc_root, &cudart_root }) { @@ -152,7 +159,7 @@ struct toolkit { } std::vector include_dirs() const { std::vector out; - for (auto const* r : { &cudart_root, &crt_root, &nvcc_root }) + for (auto const* r : { &cudart_root, &crt_root, &nvcc_root, &curand_root }) if (!r->empty() && std::filesystem::is_directory(*r + "/include")) out.push_back(*r + "/include"); return out; @@ -177,6 +184,7 @@ inline std::optional find_toolkit() { t.nvcc_root = xpkg("cuda-nvcc"); t.cudart_root = xpkg("cuda-cudart"); t.crt_root = xpkg("cuda-crt"); + t.curand_root = xpkg("libcurand"); t.driver_dir = xpkg("libcuda-host-link"); if (t.nvcc_root.empty() || t.cudart_root.empty()) { std::println(std::cerr, @@ -432,6 +440,19 @@ inline std::vector plan(std::span sources, options opt std::println(std::cerr, "mcpp.rules.cuda: the clang route needs the toolchain's clang++ at {}", driver_cc); return out; } + // Refused here rather than at clang's include error: the header it + // would fail on belongs to a payload the project has to name, and the + // diagnostic names it. A host copy is never searched for -- that is + // how the leak above survived every local build. + if (tk->curand_root.empty() + || !std::filesystem::exists(tk->curand_root + "/include/curand_mtgp32_kernel.h")) { + std::println(std::cerr, + "mcpp.rules.cuda: the clang route needs cuRAND's headers, which clang's CUDA " + "wrapper includes unconditionally.\n" + " Name the payload under [xlings.workspace] and mcpp provisions it on first use:\n" + " \"xim:libcurand\" = \"10.3.10.19\" (the 12.9 line; 10.4.x pairs with 13.x)"); + return out; + } front = { driver_cc, "-x", "cuda", "-std=c++17", "-O2", "-fPIC", "--cuda-path=" + tk->nvcc_root, "-Wno-unknown-cuda-version", // ⚠️ NVIDIA'S HEADER REFUSES libc++, AND THE REFUSAL IS diff --git a/tests/cuda-consumer/mcpp.toml b/tests/cuda-consumer/mcpp.toml index 03d9f46..a720893 100644 --- a/tests/cuda-consumer/mcpp.toml +++ b/tests/cuda-consumer/mcpp.toml @@ -38,6 +38,10 @@ cuda-runtime = "2026.09.05" [xlings.workspace] "xim:cuda-nvcc" = "12.9.86" "xim:cuda-cudart" = "12.9.79" +# cuRAND's headers: clang's CUDA wrapper includes curand_mtgp32_kernel.h for +# every device unit, so the clang route needs them even when nothing calls +# cuRAND. The 10.3.x line pairs with CUDA 12.x. +"xim:libcurand" = "10.3.10.19" "xim:libcuda-host-link" = { linux = "0.0.1" } [build] From 22e49a953531c708440fe14cd5fd8fa665583aea Mon Sep 17 00:00:00 2001 From: speak-agent Date: Sat, 5 Sep 2026 20:34:48 +0800 Subject: [PATCH 3/3] fix(rules-cuda): the clang route also names the CCCL payload curand_mtgp32_kernel.h includes from CCCL, which the 12.x toolkits ship as the separate cuda-cccl package and which the host's /usr/include had supplied silently as well. The rule adds the payload's include directory (and include/cccl for the 13.x layout) and refuses the clang route without it, naming both entries; the fixture declares it. --- rules/cuda.cppm | 29 ++++++++++++++++++++++------- tests/cuda-consumer/mcpp.toml | 3 +++ 2 files changed, 25 insertions(+), 7 deletions(-) diff --git a/rules/cuda.cppm b/rules/cuda.cppm index 534c678..8c51659 100644 --- a/rules/cuda.cppm +++ b/rules/cuda.cppm @@ -147,7 +147,10 @@ struct toolkit { // cannot compile any device unit without them. Measured 2026-09-05: on a // developer machine the header was found in the HOST's /usr/include and // the leak went unnoticed until a runner with no host CUDA refused it. - std::string curand_root; + // The same header then includes from CCCL (libcu++), which the + // 12.x toolkits ship as the separate `cuda-cccl` package, and which the + // host's /usr/include had supplied in the same way. + std::string curand_root, cccl_root; std::string nvcc() const { return nvcc_root + "/bin/nvcc"; } std::string host_config() const { for (auto const* r : { &crt_root, &nvcc_root, &cudart_root }) { @@ -159,11 +162,19 @@ struct toolkit { } std::vector include_dirs() const { std::vector out; - for (auto const* r : { &cudart_root, &crt_root, &nvcc_root, &curand_root }) + for (auto const* r : { &cudart_root, &crt_root, &nvcc_root, &cccl_root, &curand_root }) if (!r->empty() && std::filesystem::is_directory(*r + "/include")) out.push_back(*r + "/include"); + // The 13.x CCCL payload nests its tree one directory down. + if (!cccl_root.empty() && std::filesystem::is_directory(cccl_root + "/include/cccl")) + out.push_back(cccl_root + "/include/cccl"); return out; } + bool has_cccl() const { + return !cccl_root.empty() + && (std::filesystem::exists(cccl_root + "/include/nv/target") + || std::filesystem::exists(cccl_root + "/include/cccl/nv/target")); + } std::vector lib_dirs() const { std::vector out; for (auto const* r : { &cudart_root, &nvcc_root }) @@ -185,6 +196,7 @@ inline std::optional find_toolkit() { t.cudart_root = xpkg("cuda-cudart"); t.crt_root = xpkg("cuda-crt"); t.curand_root = xpkg("libcurand"); + t.cccl_root = xpkg("cuda-cccl"); t.driver_dir = xpkg("libcuda-host-link"); if (t.nvcc_root.empty() || t.cudart_root.empty()) { std::println(std::cerr, @@ -444,13 +456,16 @@ inline std::vector plan(std::span sources, options opt // would fail on belongs to a payload the project has to name, and the // diagnostic names it. A host copy is never searched for -- that is // how the leak above survived every local build. - if (tk->curand_root.empty() - || !std::filesystem::exists(tk->curand_root + "/include/curand_mtgp32_kernel.h")) { + const bool curand_ok = !tk->curand_root.empty() + && std::filesystem::exists(tk->curand_root + "/include/curand_mtgp32_kernel.h"); + if (!curand_ok || !tk->has_cccl()) { std::println(std::cerr, "mcpp.rules.cuda: the clang route needs cuRAND's headers, which clang's CUDA " - "wrapper includes unconditionally.\n" - " Name the payload under [xlings.workspace] and mcpp provisions it on first use:\n" - " \"xim:libcurand\" = \"10.3.10.19\" (the 12.9 line; 10.4.x pairs with 13.x)"); + "wrapper includes unconditionally, and CCCL's, which they include in turn.\n" + " Name the payloads under [xlings.workspace] and mcpp provisions them on first use:\n" + " \"xim:cuda-cccl\" = \"12.9.27\" (the 12.9 line; 13.x pairs with 13.x)\n" + " \"xim:libcurand\" = \"10.3.10.19\" (the 12.9 line; 10.4.x pairs with 13.x)\n" + " (found cccl: '{}', curand: '{}')", tk->cccl_root, tk->curand_root); return out; } front = { driver_cc, "-x", "cuda", "-std=c++17", "-O2", "-fPIC", diff --git a/tests/cuda-consumer/mcpp.toml b/tests/cuda-consumer/mcpp.toml index a720893..003c625 100644 --- a/tests/cuda-consumer/mcpp.toml +++ b/tests/cuda-consumer/mcpp.toml @@ -42,6 +42,9 @@ cuda-runtime = "2026.09.05" # every device unit, so the clang route needs them even when nothing calls # cuRAND. The 10.3.x line pairs with CUDA 12.x. "xim:libcurand" = "10.3.10.19" +# CCCL (libcu++, cub, thrust): cuRAND's header includes from it. +# The 12.x toolkits ship it as a separate package; 12.9.27 is the 12.9 line. +"xim:cuda-cccl" = "12.9.27" "xim:libcuda-host-link" = { linux = "0.0.1" } [build]