From 6493c958207b3946604bf2f0100742e11078b489 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Mon, 7 Sep 2026 22:18:40 +0800 Subject: [PATCH 01/20] 0.3.0: one surface for every embedded payload, and a rule for Slang A consumer of `mcpp.rules.spirv` or `mcpp.tools.embed` had to include a generated header by name. Both now hand their payloads to one generator, `mcpp::plugins::surface`, and a consumer writes `import myapp.shaders;` instead. `mcpp.rules.slang` joins them, producing the same declarations from a different language. WHAT A CONSUMER NAMES The module name and the namespace are one identifier path, and a payload's directory below the globbed tree adds a segment: `shaders/post/tone.comp` becomes `myapp::shaders::post::tone_comp()`. The name is derived from the package unless the project sets `options::module_name`. The interface is a function rather than a variable, because `constexpr` and `extern` are mutually exclusive and a variable could not keep one shape across the ways bytes can be stored. It names no standard-library type: measured with GCC 16.1 on a 1 MB payload, an interface returning `std::span` produced a 1 313 968-byte BMI against 1 808 bytes for the std-free equivalent, and that cost is fixed rather than proportional to the payload -- it is ``'s templates, present whether the payload is 16 KB or 16 MB. The accessors have C language linkage so an ordinary translation unit can define them. Declaring them with module linkage would have needed a module implementation unit, and that unit includes the data headers, which would put the arrays back into the interface's own compilation. THE DEFAULT FOLLOWS THE PROJECT, AND AN OLDER ENGINE KEEPS ITS BEHAVIOUR `[language] modules` decides the surface, reported as `MCPP_LANGUAGE_MODULES`. An engine that does not report it leaves the header surface in place, which is what every consumer of this package had before, so the upgrade needs no project to declare anything. TWO DEFECTS THE NEW FIXTURES FOUND Two shaders with one stem in two directories produced one array, silently. Their generated headers were byte-identical, and GCC's `#pragma once` treats two files with the same size and content as the same file, so the second include did nothing and both accessors returned the first array. The program printed the right magic number twice. The array's name now carries the directory, and the fixture asserts the two pointers differ rather than only that both look like SPIR-V. With no subdirectory the name is unchanged. `tools.embed`'s `_size` counts elements while the surface reports bytes, and `null_terminate` appends a byte that `_size` does not count -- so neither `_size` nor `sizeof` is the answer on its own. Items now carry the size expression their producer knows to be right. ONE PACKAGE, ONE VERSION `mcpp::plugins::version` said `0.1.1` while the package said `0.2.6`, which nothing noticed because nothing read it. Every rule now states it with `mcpp::fact`, and CI compares the two spellings. SLANG A different language, not a second driver for GLSL: its own module system, generics, and targets beyond SPIR-V, for which the Vulkan axis `rules.spirv` reads has no answer. `slangc`'s embedded output names `uint32_t` and `size_t` and includes nothing, which is the defect 0.2.6 fixed for glslang, so this rule writes the public header around it from the start. The size comes from the `_sizeInBytes` constant the compiler emits rather than from `sizeof`. A Slang module package needs no new concept: it is an ordinary mcpp package whose `include_dirs` names its `.slang` directory. `xim:slang` is pinned exactly rather than floored. Slang's version is coupled to no driver, so the rule this package follows would make it a `>=` -- but a `>=` cannot reach xlings on Windows, where cmd.exe reads the `>` as a redirection, and a floor holding on two platforms of three would resolve different payloads for one project depending on the machine that built it. `rules.slang` needs an engine whose device-source table lists `.slang`. A constrained glob's `accel` key does not make a file a device source; the engine's `kDeviceExtensions` does, and before it listed `.slang` the file fell through to the ordinary source scan and was refused. --- .github/workflows/ci.yml | 112 +++++ README.md | 47 +- mcpp.toml | 23 +- rules/slang.cppm | 454 ++++++++++++++++++ rules/spirv.cppm | 193 +++++++- src/plugins.cppm | 413 +++++++++++++++- tests/all-rules-compile/build.mcpp | 2 + tests/all-rules-compile/mcpp.toml | 3 +- tests/embed-module-consumer/build.mcpp | 20 + tests/embed-module-consumer/data/alpha.txt | 1 + tests/embed-module-consumer/data/beta.txt | 1 + tests/embed-module-consumer/mcpp.toml | 26 + tests/embed-module-consumer/src/main.cpp | 36 ++ tests/slang-consumer/build.mcpp | 15 + tests/slang-consumer/mcpp.toml | 39 ++ tests/slang-consumer/shaders/scale.slang | 29 ++ tests/slang-consumer/src/main.cpp | 25 + tests/spirv-module-consumer/build.mcpp | 15 + tests/spirv-module-consumer/mcpp.toml | 39 ++ .../shaders/a/scale.comp | 15 + .../shaders/b/scale.comp | 15 + tests/spirv-module-consumer/src/main.cpp | 44 ++ tools/embed.cppm | 84 ++++ 23 files changed, 1629 insertions(+), 22 deletions(-) create mode 100644 rules/slang.cppm create mode 100644 tests/embed-module-consumer/build.mcpp create mode 100644 tests/embed-module-consumer/data/alpha.txt create mode 100644 tests/embed-module-consumer/data/beta.txt create mode 100644 tests/embed-module-consumer/mcpp.toml create mode 100644 tests/embed-module-consumer/src/main.cpp create mode 100644 tests/slang-consumer/build.mcpp create mode 100644 tests/slang-consumer/mcpp.toml create mode 100644 tests/slang-consumer/shaders/scale.slang create mode 100644 tests/slang-consumer/src/main.cpp create mode 100644 tests/spirv-module-consumer/build.mcpp create mode 100644 tests/spirv-module-consumer/mcpp.toml create mode 100644 tests/spirv-module-consumer/shaders/a/scale.comp create mode 100644 tests/spirv-module-consumer/shaders/b/scale.comp create mode 100644 tests/spirv-module-consumer/src/main.cpp diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 0633fd8..20777a8 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -70,6 +70,22 @@ jobs: # 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. + # ONE PACKAGE, ONE VERSION. `mcpp::plugins::version` was `0.1.1` while the + # package was `0.2.6`, and nothing noticed because nothing read it. Every + # rule now states it with `mcpp::fact`, and this compares the two + # spellings so a stale constant is a red build rather than a wrong entry + # in every consumer's log. + - name: the collection states its own version + run: | + declared=$(sed -n 's/^version *= *"\(.*\)".*/\1/p' mcpp.toml | head -1) + insource=$(sed -n 's/.*version *= *"\(.*\)";.*/\1/p' src/plugins.cppm | head -1) + echo "mcpp.toml=$declared src/plugins.cppm=$insource" + if [ "$declared" != "$insource" ]; then + echo "FAIL: mcpp.toml says '$declared' and src/plugins.cppm says '$insource'" + exit 1 + fi + echo "ok: one version, two files agreeing" + - name: rules-spirv through a consumer working-directory: tests/spirv-consumer run: | @@ -77,6 +93,49 @@ jobs: "$MCPP" run | tee run.log grep -q '^magic=07230203' run.log + # THE MODULE SURFACE, AND THE COLLISION IT DISSOLVES. + # + # This fixture has `shaders/a/scale.comp` and `shaders/b/scale.comp` -- + # one stem, two directories -- which the step below this one asserts is + # REFUSED when the two would produce one output. They no longer do: the + # directory is part of the namespace, the generated file's path and the + # array's name, so both are built and reached as `shaders::a::scale_comp` + # and `shaders::b::scale_comp`. + # + # The program asserts the two arrays are DISTINCT OBJECTS, not only that + # both carry the SPIR-V magic number. Without that assertion this fixture + # passed while both accessors returned one array: the two generated data + # headers were byte-identical, and GCC's `#pragma once` treats two files + # with the same size and content as the same file, so the second include + # silently did nothing. + - name: rules-spirv through the module surface + working-directory: tests/spirv-module-consumer + run: | + "$MCPP" build + "$MCPP" run | tee run.log + grep -q '^a/scale.comp: magic=07230203' run.log + grep -q '^b/scale.comp: magic=07230203' run.log + grep -q '^all ok' run.log + # The point of the surface: a consumer names no generated file. + if grep -rn 'scale_comp\.h\|\.inc"' src/; then + echo "FAIL: a consumer source names a generated file" + exit 1 + fi + echo "ok: the consumer names only the module" + # And what it imports really is a generated module interface. + test -f target/.build-mcpp/out/spirv/spirv_module_consumer.shaders.cppm + grep -q '^export module spirv_module_consumer.shaders;' \ + target/.build-mcpp/out/spirv/spirv_module_consumer.shaders.cppm + # The interface must name no standard-library type: one that does + # carries that header's templates in every consumer's BMI, measured at + # 727 times the size of the std-free form. + if grep -n 'std::\|#include <' \ + target/.build-mcpp/out/spirv/spirv_module_consumer.shaders.cppm; then + echo "FAIL: the generated module interface names a standard-library type" + exit 1 + fi + echo "ok: the generated interface is std-free" + # 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 @@ -139,6 +198,33 @@ jobs: rm -f dup.log restored.log echo "ok: and it builds again with one shader per stem" + # SLANG: A DIFFERENT LANGUAGE, THE SAME SURFACE. + # + # The point of this step is not that Slang compiles -- it is that a + # consumer cannot tell which rule produced its shaders. The fixture's + # `main.cpp` contains nothing Slang-specific and would compile unchanged + # against `mcpp.rules.spirv` if the shader were GLSL. + # + # It needs an engine that knows the `.slang` extension: a constrained + # glob's `accel` key does not make a file a device source, the engine's + # own `kDeviceExtensions` table does, and before that table listed + # `.slang` the file fell through to the ordinary source scan and was + # refused with "mcpp has no role for the extension". + - name: rules-slang through a consumer + working-directory: tests/slang-consumer + run: | + "$MCPP" build + "$MCPP" run | tee run.log + grep -q '^scale.slang: magic=07230203' run.log + grep -q '^all ok' run.log + if grep -rn 'scale\.h\|\.inc"' src/; then + echo "FAIL: a consumer source names a generated file" + exit 1 + fi + grep -q '^export module slang_consumer.shaders;' \ + target/.build-mcpp/out/slang/slang_consumer.shaders.cppm + echo "ok: a Slang shader reached through the same surface as a GLSL one" + # 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 @@ -154,6 +240,32 @@ jobs: grep -q '^size=8 ' run2.log git checkout -- data/message.txt + # THE SAME SURFACE FROM A TOOL RATHER THAN A RULE. + # + # `file()` writes a header per input and the consumer includes each by + # name; `group()` hands the same headers to the same generator the shader + # rules use, so a set of data files and a set of shaders reach a consumer + # through one shape. + # + # The size assertion is the one that matters. `_size` counts ELEMENTS + # and the surface reports BYTES, and `null_terminate` appends a byte that + # `_size` does not count -- so neither `_size` nor `sizeof` is the answer + # on its own. The fixture compares the payload against its exact text, and + # a trailing NUL fails it. + - name: tools-embed through the module surface + working-directory: tests/embed-module-consumer + run: | + "$MCPP" build + "$MCPP" run | tee run.log + grep -q '^alpha.txt: bytes=46 ok' run.log + grep -q '^beta.txt: bytes=15 ok' run.log + grep -q '^all ok' run.log + if grep -rn 'alpha_txt\.h\|beta_txt\.h' src/; then + echo "FAIL: a consumer source names a generated file" + exit 1 + fi + echo "ok: two payloads, one import, sizes exact" + # 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. diff --git a/README.md b/README.md index 06c56af..44dab72 100644 --- a/README.md +++ b/README.md @@ -49,7 +49,8 @@ engine's own module family and is not used here. | `rules-ascendc` | `mcpp.rules.ascendc` | 2026.9.6.6 | `[build] accel = "ascend8.5+{dav-c220}"`, a constrained glob for `*.asc`. Compiles with BiSheng in MIXED mode, so the object carries the device binary and a host-callable launcher and joins the ordinary link -- no registration file and no device-link step. Its own engine needs are `.asc` in the device-source table and `mcpp::link_flag` for the `-rpath-link` the toolkit's shared libraries require, both 2026.9.6.5 | | `rules-cuda` | `mcpp.rules.cuda` | 2026.9.6.6 | `[build] accel = "cuda…"`, a constrained glob for `*.cu`; the clang route with an LLVM toolchain, the nvcc route with a GCC one | | `rules-hip` | `mcpp.rules.hip` | 2026.9.6.6 | `[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.6.6 | `[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-slang` | `mcpp.rules.slang` | see below | `[build] accel = "vulkan1.2"`, a constrained glob for `*.slang`. Slang is a different language from GLSL rather than a second driver for it -- its own module system, generics, and targets beyond SPIR-V -- so it is a rule of its own. Its engine need is `.slang` in the device-source table | +| `rules-spirv` | `mcpp.rules.spirv` | 2026.9.6.6 | `[build] accel = "vulkan1.2"`, a constrained glob for the shader stages; compiles each shader through a `role = "source"` action and states which of the two compilers produced it | | `rules-sycl` | `mcpp.rules.sycl` | 2026.9.6.6 | `[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. Its own engine need is `.sycl` in the device-source table, 2026.9.6.1 | | `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 | @@ -112,6 +113,7 @@ Each rule therefore selects the extensions it claims and leaves the rest: | `rules-ascendc` | `.asc`, `.cce` | | `rules-cuda` | `.cu` | | `rules-hip` | `.hip` | +| `rules-slang` | `.slang` | | `rules-sycl` | `.sycl` | | `rules-spirv` | `.comp .vert .frag .geom .tesc .tese .mesh .task .rgen .rint .rahit .rchit .rmiss .rcall`, and `.glsl` / `.hlsl` so that a stage-less name is refused by name rather than by absence | @@ -138,6 +140,45 @@ 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. +## What a consumer names + +A member that embeds a payload -- `rules-spirv`, `rules-slang`, `tools-embed` -- +does not leave the consumer to include a generated header. All three hand their +payloads to one generator, `mcpp::plugins::surface`, so what a consumer writes +is the same whichever produced them: + +```cpp +import myapp.shaders; + +const auto s = myapp::shaders::blur_comp(); +VkShaderModuleCreateInfo ci{ .codeSize = s.size_bytes, .pCode = s.code }; +``` + +**The module name and the namespace are one identifier path.** `myapp.shaders` +gives `myapp::shaders`, and a payload's directory below the globbed tree adds a +segment: `shaders/post/tone.comp` is `myapp::shaders::post::tone_comp()`. The +name is derived from the package unless the project sets `options::module_name`. + +**The interface is a function, and it names no standard-library type.** A +variable cannot keep one shape across the ways bytes can be stored, because +`constexpr` and `extern` are mutually exclusive. A std type in the interface is +worse than it looks: measured with GCC 16.1 on a 1 MB payload, an interface +returning `std::span` produced a 1 313 968-byte BMI against 1 808 bytes for the +std-free equivalent, and that cost is fixed rather than proportional to the +payload -- it is ``'s templates, present whether the payload is 16 KB or +16 MB. A consumer that wants a `std::span` constructs one from the two members. + +**The default follows the project.** `[language] modules = true` gives the +module surface, `false` gives a header with the same declarations. mcpp reports +the setting as `MCPP_LANGUAGE_MODULES`; an engine that does not report it leaves +the header surface in place, so an older engine keeps the behaviour every +consumer of this package had before the surface existed. + +**One copy of the bytes.** A generated data header declares a `static` array, so +before this every translation unit that included one carried its own copy. +Exactly one translation unit -- the generated implementation -- includes them +now, and every consumer reaches the same array through the accessor. + ## How the engine sees this package mcpp compiles every module interface unit among a host-module package's @@ -151,7 +192,9 @@ in the same command as the consumer's `build.mcpp`, and may import `std`, ``` mcpp.toml the package: one feature per member -src/plugins.cppm export module mcpp.plugins; +src/plugins.cppm export module mcpp.plugins; the lib root: the version, and + mcpp::plugins::surface, which every member that embeds a + payload uses to write the declarations a consumer names rules/.cppm export module mcpp.rules.; tools/.cppm export module mcpp.tools.; tests// one project per member, built by CI with the pinned mcpp diff --git a/mcpp.toml b/mcpp.toml index 18b7469..8331f9e 100644 --- a/mcpp.toml +++ b/mcpp.toml @@ -1,7 +1,7 @@ [package] name = "plugins" namespace = "mcpp" -version = "0.2.6" +version = "0.3.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"] @@ -24,6 +24,7 @@ default = [] rules-ascendc = { sources = ["rules/ascendc.cppm"] } rules-cuda = { sources = ["rules/cuda.cppm"] } rules-hip = { sources = ["rules/hip.cppm"] } +rules-slang = { sources = ["rules/slang.cppm"] } rules-spirv = { sources = ["rules/spirv.cppm"] } rules-sycl = { sources = ["rules/sycl.cppm"] } tools-embed = { sources = ["tools/embed.cppm"] } @@ -168,6 +169,26 @@ tools-embed = { sources = ["tools/embed.cppm"] } [target.'cfg(all(accelerator = "vulkan", windows))'.feature-xlings.rules-spirv] "xim:shaderc" = "2026.3" +# ONE TABLE FOR ALL THREE PLATFORMS, WHICH `rules-spirv` COULD NOT HAVE. +# +# `xim:slang` ships slangc for linux, macosx and windows and for both x86_64 and +# aarch64, so this rule does not change compilers by platform the way the GLSL +# rule does, and one table says everything. +# +# AN EXACT VERSION RATHER THAN A FLOOR, AND NOT FOR THE USUAL REASON. Slang's +# version is coupled to no driver, so the rule this package follows would make +# it a `>=`. It cannot be one on Windows: mcpp hands the provisioning request to +# xlings as JSON on a cmd.exe command line, where the `>` is read as a +# redirection (see the note above `rules-spirv`'s Windows entry). +# +# Writing `>=` on the two platforms that tolerate it and an exact version on the +# third would resolve DIFFERENT payloads for one project depending on which +# machine built it, which is how a defect comes to exist on one operating system +# only. A uniform exact version is the weaker constraint honestly applied +# everywhere; a project that wants a newer Slang names it and wins. +[target.'cfg(accelerator = "vulkan")'.feature-xlings.rules-slang] +"xim:slang" = "2026.14.1" + # 8.5.0 is a real floor rather than a preference: the mixed-mode object -- the # one carrying both the device binary and a host-callable launcher, which is # what lets it join an ordinary link -- and the 38 SoC simulators are what this diff --git a/rules/slang.cppm b/rules/slang.cppm new file mode 100644 index 0000000..0e8d0b6 --- /dev/null +++ b/rules/slang.cppm @@ -0,0 +1,454 @@ +// mcpp.rules.slang -- how a Slang translation unit becomes SPIR-V, stated once. +// +// WHY THIS IS A RULE OF ITS OWN RATHER THAN A THIRD FLAVOUR OF rules.spirv. +// +// glslang and glslc are two DRIVERS for one language: they compile the same +// `.comp`, and `mcpp.rules.spirv` chooses between them because the choice +// changes nothing a consumer sees. Slang is a different LANGUAGE. It has its +// own extension, its own module system (`import`, not `#include`), generics, +// and a target set that includes DXIL, Metal and WGSL -- for which the Vulkan +// axis `rules.spirv` reads has no answer at all. Folding it in would have put +// the rule ahead of the engine's own vocabulary. +// +// WHAT IS AND IS NOT IN SCOPE HERE. +// +// SPIR-V under a Vulkan accelerator, and nothing else. Slang can emit DXIL and +// Metal, and this rule deliberately does not, because `[build] accel` cannot +// yet express either and a rule that accepted a target the manifest could not +// name would be answering a question nobody asked it. +// +// THE COMPILER'S EMBEDDED OUTPUT NAMES TYPES IT DOES NOT INCLUDE. +// +// `-source-embed-style u32` writes a complete declaration: +// +// const uint32_t t_slang_spv[] = +// { 0x07230203, ... }; +// const size_t t_slang_spv_sizeInBytes = 172; +// +// and includes nothing, so the file names `uint32_t` and `size_t` with no +// declaration for either. That is the same defect `mcpp.rules.spirv` fixed in +// 0.2.6 for glslang's `-x --vn`, found the same way -- a program whose FIRST +// include was the generated header -- so this rule takes the same shape from +// the start: the compiler writes `.inc` and the rule writes `.h` +// around it. +// +// THE SIZE COMES FROM THE COMPILER, NOT FROM `sizeof`. +// +// `-source-embed-style` emits `_sizeInBytes` beside the array. It equals +// `sizeof` for this style and would not for a style that terminates its output, +// so the generated accessor uses what the compiler stated rather than a +// coincidence that holds today. +module; +#include +#include + +export module mcpp.rules.slang; + +import std; +import mcpp; +// The lib root, which carries `mcpp::plugins::surface` -- the declarations a +// consumer names, shared with `mcpp.rules.spirv` so a project that has both +// reaches them through one shape. +import mcpp.plugins; + +// `std::println` is avoided here for the reason every file in this package +// records: it is not header-only, and the symbols its overloads reach for were +// added to libc++ in a version macOS 14 does not ship. + +export namespace mcpp::rules::slang { + +struct options { + // The SPIR-V version, as Slang spells a profile: `spirv_1_5`. Left empty it + // is derived from the accelerator axis, which is where `[build] accel` + // already states what the build is for. + std::string profile; + // `-I` for `#include` and `import`, `-D` for the preprocessor. Relative + // entries resolve against the package root; an absolute entry passes + // through. + // + // A SLANG MODULE PACKAGE NEEDS NO NEW CONCEPT. It is an ordinary mcpp + // package whose `include_dirs` names its `.slang` directory, and a consumer + // adds that directory here with `mcpp::dep_dir`. Another build system had + // to invent a scope API for this because it had no package-level dependency + // to reuse; this one does. + std::vector includes; + std::vector defines; + // `-O` levels Slang accepts. Empty leaves the compiler's own default. + std::string optimization = "3"; + // An explicit compiler path wins over discovery. + std::string compiler; + std::string out_dir = std::string(mcpp::out_dir()); + + // ── What a consumer names ──────────────────────────────────────────────── + // Identical to `mcpp.rules.spirv`, from the same generator. See + // `mcpp::plugins::surface`. + mcpp::plugins::surface::kind surface = mcpp::plugins::surface::default_surface(); + std::string module_name; + std::string base_dir; +}; + +// Where the generated files are written. +inline std::string include_dir(const options& opt) { + return (std::filesystem::path(opt.out_dir) / "slang").string(); +} + +// ─── What the engine said ────────────────────────────────────────────────── + +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()); + auto plus = rest.find('+'); + t.version = std::string(trim(plus == std::string_view::npos ? rest + : rest.substr(0, plus))); + } + return t; +} + +// THE PROFILE IS DERIVED FROM THE VULKAN VERSION, BECAUSE THAT IS WHAT DECIDES +// IT. Each Vulkan release admits SPIR-V up to a fixed version, and a module +// emitted above it is refused by the loader at `vkCreateShaderModule` rather +// than at build time. The table is the one in the Vulkan specification's +// appendix; an unknown version falls back to the floor every Vulkan +// implementation accepts rather than guessing upward. +inline std::string profile_for(std::string_view vulkanVersion) { + if (vulkanVersion == "1.3") return "spirv_1_6"; + if (vulkanVersion == "1.2") return "spirv_1_5"; + if (vulkanVersion == "1.1") return "spirv_1_3"; + return "spirv_1_0"; +} + +// ─── The compiler ────────────────────────────────────────────────────────── + +inline bool is_file(const std::string& p) { + std::error_code ec; + return !p.empty() && std::filesystem::is_regular_file(p, ec); +} + +// Windows spells the separator differently and its paths contain the character +// the other platforms separate on, so this is not one constant used twice. +#if defined(_WIN32) +inline constexpr char kPathSep = ';'; +inline constexpr const char* kExeSuffix = ".exe"; +#else +inline constexpr char kPathSep = ':'; +inline constexpr const char* kExeSuffix = ""; +#endif + +inline std::string first_on_path(const char* exe) { + const char* path = std::getenv("PATH"); + if (!path || !*path) return {}; + std::string_view sv(path); + for (std::size_t i = 0; i <= sv.size();) { + auto sep = sv.find(kPathSep, 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) / (std::string(exe) + kExeSuffix)).string(); + if (is_file(p)) return p; + } + return {}; +} + +// Discovery, in the order a project can predict: what it named, what the +// environment named, the payload the rule declared, then the PATH. The PATH +// comes last on purpose -- a host slangc 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_SLANGC"); e && *e) return e; + if (const char* dir = mcpp::xpkg_dir("slang"); dir && *dir) + if (auto p = (std::filesystem::path(dir) / "bin" / (std::string("slangc") + kExeSuffix)).string(); + is_file(p)) return p; + return first_on_path("slangc"); +} + +inline std::string run_and_capture(const std::string& cmd) { +#if defined(_WIN32) + FILE* p = ::_popen(cmd.c_str(), "r"); +#else + FILE* p = ::popen(cmd.c_str(), "r"); +#endif + if (!p) return {}; + std::string text; + char buf[512]; + while (std::fgets(buf, sizeof buf, p)) text += buf; +#if defined(_WIN32) + ::_pclose(p); +#else + ::pclose(p); +#endif + return text; +} + +// `slangc -v` prints the release on its own, e.g. `2026.14.1`. +inline std::string compiler_version(const std::string& exe) { + auto text = run_and_capture("\"" + exe + "\" -v 2>&1"); + for (auto& c : text) if (c == '\r') c = '\n'; + auto nl = text.find('\n'); + return std::string(trim(nl == std::string::npos ? text : text.substr(0, nl))); +} + +// ─── This rule's share of the device sources ─────────────────────────────── +// +// Each rule takes the extensions it CLAIMS and leaves the rest to whoever +// claims those, for the reason `mcpp.rules.spirv` records: a project with two +// backends puts every device source in one list, and a rule that took all of it +// would hand a compiler a file it does not accept and report the file's +// contents rather than the rule that should have had it. +constexpr std::string_view kClaimed[] = { ".slang" }; + +inline bool claims_extension(std::string_view path) { + const auto slash = path.find_last_of("/\\"); + const auto name = slash == std::string_view::npos ? path : path.substr(slash + 1); + const auto dot = name.rfind('.'); + if (dot == std::string_view::npos) return false; + const auto ext = name.substr(dot); + for (auto e : kClaimed) if (e == ext) return true; + return false; +} + +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() && claims_extension(one)) out.emplace_back(one); + } + return out; +} + +// ─── Naming ──────────────────────────────────────────────────────────────── + +// `shaders/post/tone.slang` under a base of `shaders` -> `{"post"}`. The same +// derivation `mcpp.rules.spirv` uses, because a project with both must not have +// to learn two. +inline std::string common_base_dir(std::span shaders) { + std::vector prefix; + bool first = true; + for (auto const& src : shaders) { + std::vector segs; + for (auto const& part : std::filesystem::path(src).parent_path()) + if (auto s = part.string(); !s.empty() && s != ".") segs.push_back(s); + if (first) { prefix = std::move(segs); first = false; continue; } + std::size_t keep = 0; + while (keep < prefix.size() && keep < segs.size() && prefix[keep] == segs[keep]) ++keep; + prefix.resize(keep); + } + std::string out; + for (auto const& s : prefix) { if (!out.empty()) out += '/'; out += s; } + return out; +} + +inline std::vector namespace_of(std::string_view src, std::string_view base) { + std::vector out; + auto dir = std::filesystem::path(src).parent_path().string(); + if (!base.empty() && dir.size() >= base.size() && dir.compare(0, base.size(), base) == 0) + dir.erase(0, base.size()); + for (auto const& part : std::filesystem::path(dir)) { + auto s = part.string(); + if (s.empty() || s == "." || s == "/") continue; + std::string seg; + for (char c : s) + seg += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; + if (std::isdigit(static_cast(seg.front()))) seg.insert(seg.begin(), '_'); + out.push_back(std::move(seg)); + } + return out; +} + +// The array the embedded output declares. The namespace path is part of it for +// the reason `mcpp.rules.spirv` records: one translation unit includes every +// generated data header, and two byte-identical headers collapse silently under +// GCC's `#pragma once`. +inline std::string symbol_of(std::span name_space, std::string_view stem) { + std::string s; + for (auto const& seg : name_space) { s += seg; s += '_'; } + for (char c : stem) + s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; + s += "_slang_spv"; + return s; +} + +// ─── The rule ────────────────────────────────────────────────────────────── + +// The public header around the compiler's embedded output. See the note at the +// top: `-source-embed-style u32` names `uint32_t` and `size_t` and includes +// nothing, so a program whose first include is the generated header does not +// compile without this. +inline bool write_header(const std::string& header, const std::string& inc) { + std::ofstream out{header, std::ios::trunc}; + if (!out) { + std::cerr << std::format("mcpp.rules.slang: cannot write {}", header) << '\n'; + return false; + } + out << "// Generated by mcpp.rules.slang.\n" + "// slangc's embedded output names `uint32_t` and `size_t` and includes\n" + "// nothing; what this adds is the types it names and a guard.\n" + "#pragma once\n" + "#include \n" + "#include \n" + "#include \"" << std::filesystem::path(inc).filename().string() << "\"\n"; + return out.good(); +} + +inline bool compile(std::span shaders, options opt = {}) { + if (shaders.empty()) return true; + + const auto cc = find_compiler(opt); + if (cc.empty()) { + std::cerr << + "mcpp.rules.slang: no Slang compiler found.\n" + " This rule DECLARES xim:slang, so a project normally writes nothing. Check,\n" + " in order: mcpp older than 2026.9.6.6; `features = [\"rules-slang\"]` missing\n" + " from the [build-dependencies] edge; or a build that names no Vulkan\n" + " accelerator.\n" + " To pin a different version, name it in your own project and it wins:\n" + " [target.'cfg(accelerator = \"vulkan\")'.xlings.workspace]\n" + " \"xim:slang\" = \"2026.14.1\"\n" + " or name the program: MCPP_SLANGC=/path/to/slangc, or set options::compiler.\n"; + return false; + } + if (const auto v = compiler_version(cc); !v.empty()) mcpp::fact("slangc", v.c_str()); + + auto profile = opt.profile; + if (profile.empty()) profile = profile_for(parse_target(mcpp::accel()).version); + + const std::string root = mcpp::manifest_dir(); + const auto gen = include_dir(opt); + std::error_code ec; + std::filesystem::create_directories(gen, ec); + + const std::string baseDir = opt.base_dir.empty() ? common_base_dir(shaders) : opt.base_dir; + const std::string moduleName = + opt.module_name.empty() + ? mcpp::plugins::surface::module_root_from_package() + ".shaders" + : opt.module_name; + + // Two shaders whose stem and directory both match would produce one output, + // which is refused here rather than left to whichever action ran last. + { + std::map seen; + for (auto const& src : shaders) { + std::string key; + for (auto const& seg : namespace_of(src, baseDir)) key += seg + "/"; + key += std::filesystem::path(src).stem().string(); + auto [it, fresh] = seen.try_emplace(key, src); + if (!fresh) { + std::cerr << std::format( + "mcpp.rules.slang: two shaders map to one output.\n" + " {}\n {}\n" + " both produce `{}.h`. The shader's directory below `{}` is part of\n" + " the name, so this is two shaders with one name in one directory.\n" + " fix: rename one of them, or compile only one.", + it->second, src, key, + baseDir.empty() ? std::string("the package root") : baseDir) << '\n'; + return false; + } + } + } + + std::vector items; + + for (auto const& src : shaders) { + const std::filesystem::path p(src); + const auto ns = namespace_of(src, baseDir); + const auto sym = symbol_of(ns, p.stem().string()); + auto dir = std::filesystem::path(gen); + for (auto const& seg : ns) dir /= seg; + std::filesystem::create_directories(dir, ec); + const auto base = (dir / p.stem().string()).string(); + const auto header = base + ".h"; + const auto inc = base + ".inc"; + const auto input = p.is_absolute() ? src : root + "/" + src; + + if (!write_header(header, inc)) return false; + + std::string headerRel; + for (auto const& seg : ns) headerRel += seg + "/"; + headerRel += p.stem().string() + ".h"; + items.push_back({ .identifier = p.stem().string(), + .name_space = ns, + .data_header = headerRel, + .data_symbol = sym, + // What the compiler stated, rather than `sizeof`. + .data_size_expr = sym + "_sizeInBytes" }); + + const std::string id = "slang:" + src; + const std::string desc = "slangc " + src; + + mcpp::action a; + a.id = id.c_str(); + a.role = "source"; // a header: ordered before compilation + a.description = desc.c_str(); + a.arg(cc.c_str()); + a.arg("-target"); a.arg("spirv"); + a.arg("-profile"); a.arg(profile.c_str()); + if (!opt.optimization.empty()) a.arg(("-O" + opt.optimization).c_str()); + 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()); + // The embedded form, and the name it declares. + a.arg("-source-embed-style"); a.arg("u32"); + a.arg("-source-embed-name"); a.arg(sym.c_str()); + a.arg("-o"); a.arg(inc.c_str()); + a.arg(input.c_str()); + a.input(input.c_str()); + a.output(inc.c_str()); + a.submit(); + } + + mcpp::include_dir(gen.c_str()); + + mcpp::plugins::surface::options so; + so.surface = opt.surface; + so.elem = mcpp::plugins::surface::element::word32; + so.module_name = moduleName; + so.out_dir = gen; + so.produced_by = "mcpp.rules.slang"; + + const auto out = mcpp::plugins::surface::emit(items, so); + if (!out) return false; + mcpp::generated(out->interface_file.c_str()); + mcpp::generated(out->impl_file.c_str()); + if (!out->include_dir.empty()) mcpp::include_dir(out->include_dir.c_str()); + + mcpp::fact("mcpp.plugins", std::string(mcpp::plugins::version).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 every other rule in this package has. +inline bool compile(options opt = {}) { + if (!*mcpp::accel()) return true; + if (!parse_target(mcpp::accel()).present) return true; + const auto shaders = device_shaders(); + if (shaders.empty()) { + mcpp::warning("[build] accel names vulkan but no constrained glob matched a " + "`.slang`; nothing was compiled for it"); + return true; + } + return compile(std::span(shaders), std::move(opt)); +} + +} // namespace mcpp::rules::slang diff --git a/rules/spirv.cppm b/rules/spirv.cppm index 070dbdb..3823726 100644 --- a/rules/spirv.cppm +++ b/rules/spirv.cppm @@ -59,6 +59,9 @@ export module mcpp.rules.spirv; import std; import mcpp; +// The lib root, which carries `mcpp::plugins::surface` -- the declarations a +// consumer names, written once for every member that embeds a payload. +import mcpp.plugins; // WHY NOTHING HERE USES `std::println`, AND WHY THAT IS NOT A STYLE CHOICE. @@ -97,6 +100,31 @@ struct options { // pins a glslang other than the one the workspace installed. std::string compiler; std::string out_dir = std::string(mcpp::out_dir()); + + // ── What a consumer names ──────────────────────────────────────────────── + // + // The declarations are written by `mcpp.plugins.surface`, which states why + // they take the shape they do. What this rule decides is only which surface + // and under which name. + + // `module_` or `c_header`. The default follows `[language] modules`, which + // mcpp reports in `MCPP_LANGUAGE_MODULES`; an engine that does not report it + // leaves the header surface in place, so an older engine keeps the behaviour + // every consumer of this package had before the surface existed. + mcpp::plugins::surface::kind surface = mcpp::plugins::surface::default_surface(); + + // The module a consumer imports, and the namespace the declarations sit in: + // `myapp.shaders` gives `myapp::shaders::blur_comp()`. Empty derives it from + // the package directory, so a project that states nothing still gets a name + // no other package in the build can claim. + std::string module_name; + + // The directory shader paths are made relative to when deriving namespaces. + // `shaders/post/tonemap.frag` under a base of `shaders` becomes + // `myapp::shaders::post::tonemap_frag`. Empty derives the base from the + // shallowest directory every shader shares, which is what a project that + // globs one tree already means and is why this is rarely written. + std::string base_dir; }; // Where the generated headers are written, and what the build program passes @@ -379,8 +407,10 @@ inline std::string_view stage_of(std::string_view ext) { // `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) { +inline std::string symbol_of(std::span name_space, + std::string_view stem, std::string_view stage) { std::string s; + for (auto const& seg : name_space) { s += seg; s += '_'; } for (char c : stem) s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; s += '_'; @@ -389,6 +419,62 @@ inline std::string symbol_of(std::string_view stem, std::string_view stage) { return s; } +// The flat case, which is every shader that sits directly in the globbed tree. +inline std::string symbol_of(std::string_view stem, std::string_view stage) { + return symbol_of(std::span{}, stem, stage); +} + +// THE BASE DIRECTORY IS DERIVED, NOT ASKED FOR. +// +// A shader's namespace comes from where it sits relative to the tree the +// project globbed, so something has to say where that tree starts. Asking the +// project would put a second spelling of the glob in the manifest, and the two +// would disagree the first time a glob moved. The shallowest directory every +// shader shares is the same answer without the second spelling: for +// `shaders/*.comp` it is `shaders` and every namespace is empty; for +// `shaders/a/x.comp` and `shaders/b/y.comp` it is still `shaders`, and the two +// land in `::a` and `::b`. +// +// A single shader has no common prefix with anything, so its own directory is +// the base and its namespace is empty -- which is the same answer the general +// case gives once a second shader appears beside it. +inline std::string common_base_dir(std::span shaders) { + std::vector prefix; + bool first = true; + for (auto const& src : shaders) { + std::vector segs; + for (auto const& part : std::filesystem::path(src).parent_path()) + if (auto s = part.string(); !s.empty() && s != ".") segs.push_back(s); + if (first) { prefix = std::move(segs); first = false; continue; } + std::size_t keep = 0; + while (keep < prefix.size() && keep < segs.size() && prefix[keep] == segs[keep]) ++keep; + prefix.resize(keep); + } + std::string out; + for (auto const& s : prefix) { if (!out.empty()) out += '/'; out += s; } + return out; +} + +// The namespace segments a shader sits in, below the group's own: the path from +// the base directory to the shader, sanitised one segment at a time. `..` cannot +// appear, because the base is a prefix of every shader by construction. +inline std::vector namespace_of(std::string_view src, std::string_view base) { + std::vector out; + auto dir = std::filesystem::path(src).parent_path().string(); + if (!base.empty() && dir.size() >= base.size() && dir.compare(0, base.size(), base) == 0) + dir.erase(0, base.size()); + for (auto const& part : std::filesystem::path(dir)) { + auto s = part.string(); + if (s.empty() || s == "." || s == "/") continue; + std::string seg; + for (char c : s) + seg += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; + if (std::isdigit(static_cast(seg.front()))) seg.insert(seg.begin(), '_'); + out.push_back(std::move(seg)); + } + return out; +} + // 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 @@ -539,6 +625,17 @@ inline bool compile(std::span shaders, options opt = {}) { std::error_code ec; std::filesystem::create_directories(gen, ec); + // Where namespaces start counting from. Derived unless the project said, + // for the reason `common_base_dir` records. + const std::string baseDir = opt.base_dir.empty() ? common_base_dir(shaders) + : opt.base_dir; + // The module a consumer imports. `.shaders` unless the project + // named one, so two packages in one build cannot claim the same module. + const std::string moduleName = + opt.module_name.empty() + ? mcpp::plugins::surface::module_root_from_package() + ".shaders" + : opt.module_name; + // TWO SHADERS THAT DIFFER ONLY BY DIRECTORY PRODUCE ONE HEADER AND ONE // SYMBOL, AND THAT HAS TO BE REFUSED HERE. // @@ -557,30 +654,43 @@ inline bool compile(std::span shaders, options opt = {}) { // it survived: a graphics project organising shaders by purpose is the // first to have two. { - std::map seen; // output stem -> first source + std::map seen; // output path -> first source for (auto const& src : shaders) { const std::filesystem::path p(src); const auto stage = stage_of(p.extension().string()); if (stage.empty()) continue; // reported below, per source - const auto key = p.stem().string() + "_" + std::string(stage); + // THE DIRECTORY IS PART OF THE NAME NOW, SO THIS FIRES LESS OFTEN. + // + // Before the surface existed, every shader's header and symbol came + // from its stem alone, so `a/scale.comp` and `b/scale.comp` collided + // and had to be refused. Both now land in their own namespace and + // their own subdirectory of the generated tree, so the refusal is + // for what it was always about: two shaders that are genuinely the + // same name in the same place. + std::string key; + for (auto const& seg : namespace_of(src, baseDir)) key += seg + "/"; + key += p.stem().string() + "_" + std::string(stage); auto [it, fresh] = seen.try_emplace(key, src); if (!fresh) { std::cerr << std::format("mcpp.rules.spirv: two shaders map to one output.\n" " {}\n" " {}\n" - " both produce `{}.h` declaring `{}`, because the name is the " - "shader's stem\n" - " and its stage -- the directory is not part of it, and could not " - "be: two\n" - " headers reaching one translation unit would still collide on the " - "symbol.\n" + " both produce `{}.h` declaring `{}`. The shader's directory below\n" + " `{}` is part of the name, so this is two shaders with one name in\n" + " one directory rather than two directories sharing a stem.\n" " fix: rename one of them, or compile only one.", - it->second, src, key, symbol_of(p.stem().string(), stage)) << '\n'; + it->second, src, key, symbol_of(p.stem().string(), stage), + baseDir.empty() ? std::string("the package root") : baseDir) << '\n'; return false; } } } + // Collected while the actions are submitted and handed to + // `mcpp.plugins.surface` afterwards, so the declarations a consumer reads + // are written once for the whole group rather than once per shader. + std::vector items; + for (auto const& src : shaders) { const std::filesystem::path p(src); const auto stage = stage_of(p.extension().string()); @@ -590,10 +700,40 @@ inline bool compile(std::span shaders, options opt = {}) { ".tese .mesh .task .rgen .rint .rahit .rchit .rmiss .rcall", src) << '\n'; return false; } - const auto sym = symbol_of(p.stem().string(), stage); - const auto base = (std::filesystem::path(gen) - / (p.stem().string() + "_" + std::string(stage))).string(); + // The generated tree mirrors the shader tree below the base directory, + // so two shaders sharing a stem in different directories produce + // different files as well as different namespaces. + const auto ns = namespace_of(src, baseDir); + // THE ARRAY'S NAME CARRIES THE DIRECTORY, AND IT HAS TO. + // + // One translation unit includes every generated data header, so two + // shaders sharing a stem would declare one name twice. That is not + // caught as a redefinition, which is what makes it worth a comment: + // GCC's `#pragma once` treats two files with the same size and the same + // content as the same file, so two identical headers -- which is exactly + // what the same shader in two directories produces -- SILENTLY collapse + // to one, and both accessors return the same array. Measured on a + // fixture with `shaders/a/scale.comp` and `shaders/b/scale.comp`: the + // program printed the right magic number twice and the two pointers + // were equal. + // + // With no subdirectory the namespace is empty and the name is what it + // has always been, so nothing an existing project generated changes. + const auto sym = symbol_of(ns, p.stem().string(), stage); + auto dir = std::filesystem::path(gen); + for (auto const& seg : ns) dir /= seg; + std::filesystem::create_directories(dir, ec); + const auto base = (dir / (p.stem().string() + "_" + std::string(stage))).string(); const auto header = base + ".h"; + // As an `#include` writes it: relative to `gen`, which is the directory + // this rule puts on the include path. + std::string headerRel; + for (auto const& seg : ns) headerRel += seg + "/"; + headerRel += p.stem().string() + "_" + std::string(stage) + ".h"; + items.push_back({ .identifier = p.stem().string() + "_" + std::string(stage), + .name_space = ns, + .data_header = headerRel, + .data_symbol = sym }); const auto input = std::filesystem::path(src).is_absolute() ? src : root + "/" + src; @@ -654,7 +794,34 @@ inline bool compile(std::span shaders, options opt = {}) { a.submit(); } + // The include path carries the generated data headers so the generated + // implementation can reach them. It is not how a consumer reaches a shader: + // that is the surface below, and no consumer writes one of these names. mcpp::include_dir(gen.c_str()); + + // ── What a consumer names ──────────────────────────────────────────────── + mcpp::plugins::surface::options so; + so.surface = opt.surface; + so.elem = mcpp::plugins::surface::element::word32; + so.module_name = moduleName; + so.out_dir = gen; + so.produced_by = "mcpp.rules.spirv"; + + const auto out = mcpp::plugins::surface::emit(items, so); + if (!out) return false; + + // Both generated files are written above, so the ordinary source scan sees + // real content rather than a placeholder -- which is what lets a generated + // module interface be an ordinary node in the module graph with nothing + // declared about it. + mcpp::generated(out->interface_file.c_str()); + mcpp::generated(out->impl_file.c_str()); + if (!out->include_dir.empty()) mcpp::include_dir(out->include_dir.c_str()); + + // Which collection produced these actions. The version was `0.1.1` while + // the package was `0.2.6` for as long as nothing read it; a fact is a + // reader, and a stale constant is now visible in every build log. + mcpp::fact("mcpp.plugins", std::string(mcpp::plugins::version).c_str()); return true; } diff --git a/src/plugins.cppm b/src/plugins.cppm index dbb427f..7ca8f0f 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -1,16 +1,419 @@ -// mcpp.plugins: the identity unit of the collection. +// mcpp.plugins: the collection's shared library. // // 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. +// request names it. This unit is the lib root: it is compiled before every +// member, so a member may import it. +// +// IT HOLDS WHAT MORE THAN ONE MEMBER NEEDS, AND ONLY THAT. +// +// The lib root is the only unit every member can import. A second unit beside +// it in `[build] sources` is NOT compiled as a host module ahead of the +// members. Measured: a member importing a second lib-root unit failed with +// +// mcpp.plugins.surface: error: failed to read compiled module +// note: imports must be built before being imported +// +// because only the lib root is built first. So shared code lives here rather +// than in a file of its own, and a member that needs it writes +// `import mcpp.plugins;`. +module; +#include + export module mcpp.plugins; import std; +import mcpp; export namespace mcpp::plugins { -inline constexpr std::string_view version = "0.1.1"; +// THIS MUST EQUAL `[package] version` IN mcpp.toml, AND CI CHECKS THAT IT DOES. +// +// It was `0.1.1` while the package was `0.2.6`, which nothing noticed because +// nothing read it. A value that is recorded and never read cannot be wrong in a +// way anyone sees, so the fix is not only to correct it: every rule now states +// it with `mcpp::fact`, which makes a build log answer "which collection +// produced these actions" and makes a stale constant a visible defect rather +// than a dormant one. +// +// One package, one version: the number lives in mcpp.toml, and the CI step +// `the collection states its own version` compares the two. +inline constexpr std::string_view version = "0.3.0"; } // namespace mcpp::plugins + +// mcpp.plugins.surface -- the interface a consumer names for an embedded payload. +// +// WHY THIS IS SHARED RATHER THAN PER-RULE. +// +// `mcpp.rules.spirv` and `mcpp.rules.slang` produce the same kind of thing: a +// block of bytes the program hands to a device API. `mcpp.tools.embed` produces +// it from a file that was already there. The three differ in who produces the +// bytes and disagree about nothing else, so the declaration a consumer reads is +// written once, here, and the shape is identical whichever produced it. A rule +// that wrote its own would be a second copy of a decision, and the two would +// drift the way the two shader compilers' headers drifted before 0.2.6. +// +// THE INTERFACE IS A FUNCTION, NOT A VARIABLE. +// +// A variable cannot keep one shape across the ways bytes can be stored, because +// `constexpr` and `extern` are mutually exclusive: an array compiled into a +// translation unit can be constant-evaluated and one living in a section cannot. +// A function can, so where the bytes live stays an option a project revises +// without touching a consumer. +// +// THE INTERFACE NAMES NO STANDARD-LIBRARY TYPE, AND THAT IS MEASURED. +// +// The same 1 MB payload, reached four ways, compiled with GCC 16.1: +// +// data in the module (`export inline constexpr`) BMI 6 282 240 B +// data in a header source 3 145 728 B +// data in an object, interface returns std::span BMI 1 313 968 B +// data in an object, interface returns a std-free POD BMI 1 808 B +// +// The third row is 727 times the fourth, and its cost is FIXED rather than +// proportional to the payload: the 1.28 MB is ``'s templates, present +// whether the payload is 16 KB or 16 MB. A consumer that wants a `std::span` +// constructs one from the two members, and `` is then included by the +// consumer that uses it rather than by every consumer that imports this. +// +// There is a second, independent reason for the same decision. `import std;` +// requires `std.gcm` to have been built, which a project with +// `modules = true, import_std = false` has not done. A std-free interface needs +// neither. +// +// WHY THE ACCESSORS HAVE C LANGUAGE LINKAGE. +// +// A function declared in a module interface has module linkage and can only be +// defined by a unit attached to that module. Defining them would therefore need +// a module implementation unit, and the definitions include the generated data +// headers -- which would put the arrays back into the interface's own +// compilation. Declaring the accessors `extern "C"` instead gives them external +// linkage, so an ordinary translation unit defines them, the module interface +// holds two declarations and one inline call per payload, and the bytes are +// never seen by the interface at all. +// +// It also removes a portability question: module implementation units are the +// least exercised corner of every implementation this package supports, and +// nothing here needs them. +// +// ONE COPY OF THE BYTES. +// +// A generated data header declares `static const uint32_t []`, so every +// translation unit that includes it gets its own copy. Under this surface +// exactly one translation unit includes it -- the generated implementation -- +// and every consumer reaches the same array through the accessor. +export namespace mcpp::plugins::surface { + +// How a consumer names the payloads. +// +// `module_` is the default where the project builds C++ modules, and +// `c_header` is what a project without them gets. The choice does not change +// any declaration's shape: the same struct, the same function names, the same +// namespaces. Only the file a consumer reaches them through differs. +enum class kind { module_, c_header }; + +// The element the accessor hands back. `word32` is what an API taking +// `const uint32_t*` wants -- `vkCreateShaderModule` is the case this package +// has -- and asking for it here is cheaper than a reinterpret_cast at every +// call site, which is undefined behaviour on an under-aligned byte array. +enum class element { byte_, word32 }; + +// One payload in a group. +struct item { + // The C++ identifier, already sanitised by the caller: `blur_comp`. + std::string identifier; + // Namespace segments BELOW the group's own, from the payload's directory + // relative to its root: `{"post"}` for `shaders/post/tonemap.frag`. This is + // what keeps two payloads with one stem from colliding, and what makes the + // name a consumer writes say where the payload came from. + std::vector name_space; + // The generated data header this payload's bytes arrive in, AS AN + // `#include` WRITES IT -- relative to the directory the caller put on the + // include path, not an absolute path and not a bare file name. The + // generated tree mirrors the payload tree, so two payloads sharing a stem + // in different directories reach different headers, and a bare file name + // could not tell them apart. + std::string data_header; + // The array that header declares, qualified if it sits in a namespace. + std::string data_symbol; + // The payload's size IN BYTES, as an expression valid where the generated + // implementation writes it. Empty means `sizeof `. + // + // IT EXISTS BECAUSE `sizeof` IS NOT ALWAYS THE ANSWER. `mcpp.tools.embed` + // can append a terminating zero byte that its own `_size` constant does not + // count, so `sizeof` and `_size` differ by one and only one of them is what + // the payload is. Measured: a group of null-terminated text payloads + // reported one byte too many and the fixture's comparison failed on a + // trailing NUL. A caller that knows the difference states it here rather + // than letting the generator guess. + std::string data_size_expr; +}; + +struct options { + kind surface = kind::module_; + element elem = element::word32; + // The module a consumer imports, e.g. `myapp.shaders`. The namespace is + // this name with `.` replaced by `::`, which is why there is no second + // field for it: two spellings of one identity drift. + std::string module_name; + // Where the generated interface and implementation are written. The caller + // adds this to the include path when the surface is `c_header`. + std::string out_dir; + // A short phrase naming what produced the payloads, for the generated + // files' first line: "mcpp.rules.spirv". + std::string produced_by; +}; + +// What the caller hands back to mcpp. +struct emitted { + // The interface unit: a `.cppm` under `module_`, a `.h` under `c_header`. + std::string interface_file; + // The translation unit defining the accessors. Always a plain `.cpp`. + std::string impl_file; + // Non-empty under `module_`: what the interface unit provides, which the + // caller declares with `mcpp::action::provides` or lets the scan find. + std::string module_name; + // Non-empty under `c_header`: the directory to put on the include path. + std::string include_dir; +}; + +// ---- internals ------------------------------------------------------------- + +inline std::vector split_module_name(std::string_view name) { + std::vector out; + for (std::size_t i = 0; i <= name.size();) { + auto dot = name.find('.', i); + auto one = dot == std::string_view::npos ? name.substr(i) : name.substr(i, dot - i); + if (!one.empty()) out.emplace_back(one); + if (dot == std::string_view::npos) break; + i = dot + 1; + } + return out; +} + +// The linker symbol an accessor carries. It is derived from the module name and +// the payload's full namespace path rather than from the identifier alone, +// because two groups in one link unit -- a project with shaders and with +// embedded data -- would otherwise define the same symbol twice and the second +// definition would be the one nobody expected. +inline std::string accessor_base(const options& opt, const item& it) { + std::string s = "mcpp_embed"; + auto add = [&](std::string_view part) { + s += '_'; + for (char c : part) + s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; + }; + for (auto const& seg : split_module_name(opt.module_name)) add(seg); + for (auto const& seg : it.name_space) add(seg); + add(it.identifier); + return s; +} + +inline const char* element_type(element e) { + return e == element::word32 ? "unsigned int" : "unsigned char"; +} + +// `unsigned int` is 32 bits on every target this package supports, and the +// generated implementation asserts it rather than assuming it: a target where +// it is not would otherwise hand a Vulkan driver a pointer into misread data, +// and the failure would be a device error naming nothing. +inline const char* element_width_assertion(element e) { + return e == element::word32 + ? "static_assert(sizeof(unsigned int) == 4,\n" + " \"mcpp.plugins.surface: element::word32 assumes a 32-bit `unsigned int`\");\n" + : ""; +} + +inline bool write_if_different(const std::string& path, std::string_view text) { + std::error_code ec; + std::filesystem::create_directories(std::filesystem::path(path).parent_path(), ec); + if (std::ifstream in(path, std::ios::binary); in) { + std::string old((std::istreambuf_iterator(in)), std::istreambuf_iterator()); + if (old == text) return true; + } + std::ofstream out(path, std::ios::binary | std::ios::trunc); + if (!out) return false; + out.write(text.data(), static_cast(text.size())); + return static_cast(out); +} + +// Open and close the namespaces an item sits in, below the group's own. +inline std::string open_namespaces(const std::vector& segs) { + std::string s; + for (auto const& one : segs) s += "namespace " + one + " {\n"; + return s; +} +inline std::string close_namespaces(const std::vector& segs) { + std::string s; + for (auto i = segs.rbegin(); i != segs.rend(); ++i) s += "} // namespace " + *i + "\n"; + return s; +} + +// The declarations shared by both surfaces. What differs between a module +// interface and a header is the preamble and how the group's namespace is +// opened; the body below is byte-identical, which is what makes switching +// surfaces a change no consumer can observe. +// The accessor declarations, at GLOBAL scope. +// +// C language linkage makes a name the same entity whatever namespace declares +// it, so these would resolve from inside the group's namespace as well. They are +// written outside it because the generated implementation defines them at global +// scope, and a reader comparing the two files should not have to know that rule +// to see that they match. +inline std::string extern_c_declarations(std::span items, const options& opt) { + const char* elem = element_type(opt.elem); + std::string s = "extern \"C\" {\n"; + for (auto const& it : items) { + const auto base = accessor_base(opt, it); + s += std::format("const {}* {}_data();\n", elem, base); + s += std::format("unsigned long {}_size();\n", base); + } + s += "}\n"; + return s; +} + +inline std::string declarations(std::span items, const options& opt) { + const char* elem = element_type(opt.elem); + std::string s; + + s += std::format( + "// The payload, as the device API wants it. Free of standard-library\n" + "// types on purpose: see mcpp.plugins.surface.\n" + "struct payload {{\n" + " const {}* code; // e.g. VkShaderModuleCreateInfo::pCode\n" + " unsigned long size_bytes; // e.g. VkShaderModuleCreateInfo::codeSize\n" + "}};\n\n", elem); + + // Grouped by namespace path so a directory's payloads are emitted together + // and each namespace is opened once. + std::vector openNow; + auto reopen = [&](const std::vector& want) { + if (openNow == want) return; + s += close_namespaces(openNow); + s += open_namespaces(want); + openNow = want; + }; + for (auto const& it : items) { + reopen(it.name_space); + const auto base = accessor_base(opt, it); + s += std::format("inline payload {}() {{ return {{ {}_data(), {}_size() }}; }}\n", + it.identifier, base, base); + } + reopen({}); + return s; +} + +// ---- the tool --------------------------------------------------------------- + +// Writes the interface and the implementation, and returns what the caller has +// to tell mcpp about them. Nothing here reads a payload's bytes, so it runs at +// plan time and works equally for a payload the graph has not produced yet. +inline std::optional emit(std::span items, const options& opt) { + if (items.empty()) return emitted{}; + if (opt.module_name.empty()) { + std::cerr << "mcpp.plugins.surface: options::module_name is required; it names both " + "the module a consumer imports and the namespace the declarations sit in\n"; + return std::nullopt; + } + if (opt.out_dir.empty()) { + std::cerr << "mcpp.plugins.surface: options::out_dir is required\n"; + return std::nullopt; + } + + const auto segs = split_module_name(opt.module_name); + const auto by = opt.produced_by.empty() ? std::string("mcpp.plugins.surface") + : opt.produced_by; + const auto dir = std::filesystem::path(opt.out_dir); + const auto body = declarations(items, opt); + + emitted out; + + // ---- interface ---- + std::string iface; + iface += std::format("// Generated by mcpp.plugins.surface for {}. Do not edit.\n", by); + if (opt.surface == kind::module_) iface += std::format("export module {};\n\n", opt.module_name); + else iface += "#pragma once\n\n"; + iface += extern_c_declarations(items, opt); + iface += "\n"; + // `export` on the opening namespace exports everything the block contains, + // including the nested namespaces a payload's directory produced. + if (opt.surface == kind::module_) iface += "export "; + iface += open_namespaces(segs); + iface += "\n" + body + "\n"; + iface += close_namespaces(segs); + + out.interface_file = (dir / (opt.module_name + + (opt.surface == kind::module_ ? ".cppm" : ".h"))).string(); + if (!write_if_different(out.interface_file, iface)) { + std::cerr << std::format("mcpp.plugins.surface: cannot write {}\n", out.interface_file); + return std::nullopt; + } + if (opt.surface == kind::module_) out.module_name = opt.module_name; + else out.include_dir = dir.string(); + + // ---- implementation ---- + // + // The one translation unit that includes the data headers. It is a plain + // `.cpp` under both surfaces, because the accessors have C language linkage + // and so need no attachment to the module. + std::string impl; + impl += std::format("// Generated by mcpp.plugins.surface for {}. Do not edit.\n" + "//\n" + "// The only translation unit that includes the generated data headers.\n" + "// Each declares a `static` array, so this is also the only copy of the\n" + "// bytes in the program.\n", by); + for (auto const& it : items) impl += std::format("#include \"{}\"\n", it.data_header); + impl += "\n"; + impl += element_width_assertion(opt.elem); + impl += "\n"; + const char* elem = element_type(opt.elem); + for (auto const& it : items) { + const auto base = accessor_base(opt, it); + const auto size = it.data_size_expr.empty() + ? std::format("sizeof {}", it.data_symbol) + : it.data_size_expr; + impl += std::format( + "extern \"C\" const {0}* {1}_data() {{ return {2}; }}\n" + "extern \"C\" unsigned long {1}_size() {{ return {3}; }}\n", + elem, base, it.data_symbol, size); + } + + out.impl_file = (dir / (opt.module_name + ".impl.cpp")).string(); + if (!write_if_different(out.impl_file, impl)) { + std::cerr << std::format("mcpp.plugins.surface: cannot write {}\n", out.impl_file); + return std::nullopt; + } + return out; +} + +// The surface a project gets when it asks for nothing. +// +// `MCPP_LANGUAGE_MODULES` is set by mcpp from `[language] modules`. An engine +// that does not set it leaves the variable absent, and the fallback is the +// header surface -- which is what every consumer of this package had before +// this module existed. An older engine therefore keeps its behaviour and a +// newer one gets the module surface without any project asking, which is the +// whole of the upgrade path. +inline kind default_surface() { + const char* v = std::getenv("MCPP_LANGUAGE_MODULES"); + return (v && (*v == '1' || *v == 't' || *v == 'T')) ? kind::module_ : kind::c_header; +} + +// `example:my-app` -> `my_app`. The module a rule names by default is this +// followed by the group's own segment, so two packages in one build cannot +// claim the same module. +inline std::string module_root_from_package() { + std::string s; + // MCPP_MANIFEST_DIR's leaf is the package directory, which is the closest + // thing a build program is told about its own name. A rule that knows + // better passes `options::module_name` and this is not consulted. + const auto leaf = std::filesystem::path(mcpp::manifest_dir()).filename().string(); + for (char c : leaf) + s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; + if (s.empty()) s = "app"; + if (std::isdigit(static_cast(s.front()))) s.insert(s.begin(), '_'); + return s; +} + +} // namespace mcpp::plugins::surface diff --git a/tests/all-rules-compile/build.mcpp b/tests/all-rules-compile/build.mcpp index f265633..b838d9e 100644 --- a/tests/all-rules-compile/build.mcpp +++ b/tests/all-rules-compile/build.mcpp @@ -8,6 +8,7 @@ import mcpp; import mcpp.rules.ascendc; import mcpp.rules.cuda; import mcpp.rules.hip; +import mcpp.rules.slang; import mcpp.rules.spirv; import mcpp.rules.sycl; import mcpp.tools.embed; @@ -19,6 +20,7 @@ int main() { bool ok = mcpp::rules::ascendc::compile() && mcpp::rules::cuda::compile() && mcpp::rules::hip::compile() + && mcpp::rules::slang::compile() && mcpp::rules::spirv::compile() && mcpp::rules::sycl::compile(); // `tools::embed` has no accelerator gate, so it is asked the one question diff --git a/tests/all-rules-compile/mcpp.toml b/tests/all-rules-compile/mcpp.toml index 4c3f3e4..2979774 100644 --- a/tests/all-rules-compile/mcpp.toml +++ b/tests/all-rules-compile/mcpp.toml @@ -36,7 +36,8 @@ import_std = true # whose host-dependent code is compiled on one platform only. [build-dependencies.mcpp] plugins = { path = "../..", features = [ - "rules-ascendc", "rules-cuda", "rules-hip", "rules-spirv", "rules-sycl", + "rules-ascendc", "rules-cuda", "rules-hip", "rules-slang", "rules-spirv", + "rules-sycl", "tools-embed", ], host-module = true } diff --git a/tests/embed-module-consumer/build.mcpp b/tests/embed-module-consumer/build.mcpp new file mode 100644 index 0000000..7890a60 --- /dev/null +++ b/tests/embed-module-consumer/build.mcpp @@ -0,0 +1,20 @@ +import std; +import mcpp; +import mcpp.plugins; +import mcpp.tools.embed; + +int main() { + mcpp::rerun_if_changed_glob("data/*.txt"); + + mcpp::tools::embed::options opt; + // Set on purpose: the group's symbols are qualified by it, and a group() + // that passed the bare name would compile for a caller that left this empty + // and fail only here. + opt.name_space = "fixture"; + opt.null_terminate = true; + + const std::vector inputs{ "data/alpha.txt", "data/beta.txt" }; + return mcpp::tools::embed::group(inputs, "embed_module_consumer.assets", + mcpp::plugins::surface::kind::module_, opt) + ? 0 : 1; +} diff --git a/tests/embed-module-consumer/data/alpha.txt b/tests/embed-module-consumer/data/alpha.txt new file mode 100644 index 0000000..5c55be7 --- /dev/null +++ b/tests/embed-module-consumer/data/alpha.txt @@ -0,0 +1 @@ +mcpp.tools.embed group fixture, first payload diff --git a/tests/embed-module-consumer/data/beta.txt b/tests/embed-module-consumer/data/beta.txt new file mode 100644 index 0000000..34702c3 --- /dev/null +++ b/tests/embed-module-consumer/data/beta.txt @@ -0,0 +1 @@ +and the second diff --git a/tests/embed-module-consumer/mcpp.toml b/tests/embed-module-consumer/mcpp.toml new file mode 100644 index 0000000..705bb82 --- /dev/null +++ b/tests/embed-module-consumer/mcpp.toml @@ -0,0 +1,26 @@ +# The consumer that tests `mcpp.tools.embed::group()`. +# +# `embed-consumer` covers `file()`: one payload, one generated header, included +# by name. This covers the surface -- several payloads reached through one +# `import`, with no generated file named anywhere in the consumer's sources. +# +# It exists as a separate fixture rather than as more assertions on the first +# one because the two entry points have different contracts, and a fixture that +# exercised both would not say which one broke. +[package] +name = "embed-module-consumer" +namespace = "example" +version = "0.1.0" +description = "Data files reached through a generated C++ module" + +[language] +standard = "c++23" +modules = true +import_std = true + +[build-dependencies.mcpp] +plugins = { path = "../..", features = ["tools-embed"], host-module = true } + +[targets.embed-module-consumer] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/embed-module-consumer/src/main.cpp b/tests/embed-module-consumer/src/main.cpp new file mode 100644 index 0000000..de5238a --- /dev/null +++ b/tests/embed-module-consumer/src/main.cpp @@ -0,0 +1,36 @@ +// Nothing here names a generated file. Two payloads arrive through one import. +import std; +import embed_module_consumer.assets; + +namespace assets = embed_module_consumer::assets; + +namespace { + +bool check(std::string_view label, assets::payload p, std::string_view want) { + const std::string_view got{ reinterpret_cast(p.code), p.size_bytes }; + const bool ok = got == want; + std::cout << std::format("{}: bytes={} {}\n", label, p.size_bytes, + ok ? "ok" : "BAD"); + if (!ok) std::cout << std::format(" want {:?}\n got {:?}\n", want, got); + return ok; +} + +} // namespace + +int main() { + bool ok = true; + ok &= check("alpha.txt", assets::alpha_txt(), + "mcpp.tools.embed group fixture, first payload\n"); + ok &= check("beta.txt", assets::beta_txt(), "and the second\n"); + + // Two payloads must be two objects. The shader fixture found this the hard + // way: two byte-identical generated headers collapse under GCC's + // `#pragma once`, and both accessors then return one array. These two + // payloads differ, so the check is cheap insurance rather than a repeat. + if (assets::alpha_txt().code == assets::beta_txt().code) { + std::cout << "BAD: both payloads resolve to one array\n"; + ok = false; + } + std::cout << (ok ? "all ok\n" : "FAILED\n"); + return ok ? 0 : 1; +} diff --git a/tests/slang-consumer/build.mcpp b/tests/slang-consumer/build.mcpp new file mode 100644 index 0000000..89e28ad --- /dev/null +++ b/tests/slang-consumer/build.mcpp @@ -0,0 +1,15 @@ +import std; +import mcpp; +import mcpp.plugins; +import mcpp.rules.slang; + +int main() { + mcpp::rerun_if_changed_glob("shaders/**/*.slang"); + mcpp::rules::slang::options opt; + // Asked for explicitly rather than left to the default, for the reason the + // GLSL module fixture records: the default follows `MCPP_LANGUAGE_MODULES`, + // and against an engine that does not report it this fixture would silently + // get the header surface, still build, and test nothing it exists for. + opt.surface = mcpp::plugins::surface::kind::module_; + return mcpp::rules::slang::compile(opt) ? 0 : 1; +} diff --git a/tests/slang-consumer/mcpp.toml b/tests/slang-consumer/mcpp.toml new file mode 100644 index 0000000..522b872 --- /dev/null +++ b/tests/slang-consumer/mcpp.toml @@ -0,0 +1,39 @@ +# The consumer that tests mcpp.rules.slang. +# +# Slang is a different language from GLSL, not a second driver for it, which is +# why it is a rule of its own. What this fixture asserts is that the two rules +# are nevertheless indistinguishable from a consumer's side: the same generated +# module shape, the same namespace derivation, the same `payload` struct. A +# project that moved a shader from GLSL to Slang would change the file and +# nothing else. +# +# No Vulkan runtime is involved, for the reason the GLSL fixtures record: the +# property that identifies a SPIR-V module without executing it is the magic +# number in its first word. +[package] +name = "slang-consumer" +namespace = "example" +version = "0.1.0" +description = "A Slang compute shader compiled to SPIR-V through mcpp.rules.slang" + +[language] +standard = "c++23" +modules = true +import_std = true + +[build-dependencies.mcpp] +plugins = { path = "../..", features = ["rules-slang"], host-module = true } + +# NO [xlings.workspace]. `xim:slang` is declared by the rule, under the feature +# that selects it and the accelerator it serves. + +[build] +accel = "vulkan1.2" +sources = [ + "src/*.cpp", + { glob = "shaders/*.slang", accel = "vulkan1.2" }, +] + +[targets.slang-consumer] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/slang-consumer/shaders/scale.slang b/tests/slang-consumer/shaders/scale.slang new file mode 100644 index 0000000..97052f7 --- /dev/null +++ b/tests/slang-consumer/shaders/scale.slang @@ -0,0 +1,29 @@ +// The same computation the GLSL fixture runs, written in Slang: out = a*x + y. +// +// Written with a generic so the file exercises what Slang has and GLSL does not. +// One storage buffer holds all three vectors, so the host side needs one +// allocation and one descriptor. +struct Push { + float a; + uint n; +}; + +[[vk::push_constant]] +ConstantBuffer push; + +[[vk::binding(0, 0)]] +RWStructuredBuffer data; + +// A generic the compiler specialises. GLSL would need a macro and a second copy +// of the body for every type; this is the difference the language buys. +T saxpy(T a, T x, T y) { + return a * x + y; +} + +[shader("compute")] +[numthreads(64, 1, 1)] +void computeMain(uint3 tid : SV_DispatchThreadID) { + const uint i = tid.x; + if (i >= push.n) return; + data[2 * push.n + i] = saxpy(push.a, data[i], data[push.n + i]); +} diff --git a/tests/slang-consumer/src/main.cpp b/tests/slang-consumer/src/main.cpp new file mode 100644 index 0000000..b717ac9 --- /dev/null +++ b/tests/slang-consumer/src/main.cpp @@ -0,0 +1,25 @@ +// Nothing here names a generated file, and nothing here is Slang-specific: the +// same source would compile against `mcpp.rules.spirv` if the shader were GLSL. +import std; +import slang_consumer.shaders; + +namespace shaders = slang_consumer::shaders; + +int main() { + const auto s = shaders::scale(); + const bool ok = s.size_bytes >= 4 && s.code[0] == 0x07230203u; + + std::cout << std::format("scale.slang: magic={:08x} bytes={} {}\n", + s.size_bytes >= 4 ? s.code[0] : 0u, s.size_bytes, + ok ? "ok" : "BAD"); + + // The size the compiler stated, not `sizeof` of whatever it wrote. SPIR-V + // is a sequence of 32-bit words, so a size that is not a multiple of four + // means the accessor is reporting something other than the module. + if (s.size_bytes % 4 != 0) { + std::cout << "BAD: size is not a whole number of SPIR-V words\n"; + return 1; + } + std::cout << (ok ? "all ok\n" : "FAILED\n"); + return ok ? 0 : 1; +} diff --git a/tests/spirv-module-consumer/build.mcpp b/tests/spirv-module-consumer/build.mcpp new file mode 100644 index 0000000..b25f8cf --- /dev/null +++ b/tests/spirv-module-consumer/build.mcpp @@ -0,0 +1,15 @@ +import std; +import mcpp; +import mcpp.plugins; +import mcpp.rules.spirv; + +int main() { + mcpp::rerun_if_changed_glob("shaders/**/*.comp"); + mcpp::rules::spirv::options opt; + // Asked for explicitly rather than left to the default, because the default + // follows `MCPP_LANGUAGE_MODULES` and an engine older than the one that + // reports it would silently give this fixture the header surface -- which + // would still build, and would test nothing this fixture exists for. + opt.surface = mcpp::plugins::surface::kind::module_; + return mcpp::rules::spirv::compile(opt) ? 0 : 1; +} diff --git a/tests/spirv-module-consumer/mcpp.toml b/tests/spirv-module-consumer/mcpp.toml new file mode 100644 index 0000000..c3fcdf7 --- /dev/null +++ b/tests/spirv-module-consumer/mcpp.toml @@ -0,0 +1,39 @@ +# The consumer that tests the MODULE surface of mcpp.rules.spirv. +# +# What it asserts that `spirv-consumer` cannot: +# +# - a consumer reaches its shaders with `import`, and names no generated +# header anywhere in its own sources; +# - two shaders with one stem in two directories are no longer a collision: +# `shaders/a/scale.comp` and `shaders/b/scale.comp` land in +# `spirv_module_consumer::shaders::a` and `::b`; +# - the module name and the namespace are the same identifier path, derived +# from the package with nothing declared for it. +# +# No Vulkan runtime is involved, for the reason `spirv-consumer` records: the +# property that identifies a SPIR-V module without executing it is the magic +# number in its first word. +[package] +name = "spirv-module-consumer" +namespace = "example" +version = "0.1.0" +description = "GLSL compute shaders reached through a generated C++ module" + +[language] +standard = "c++23" +modules = true +import_std = true + +[build-dependencies.mcpp] +plugins = { path = "../..", features = ["rules-spirv"], host-module = true } + +[build] +accel = "vulkan1.2" +sources = [ + "src/*.cpp", + { glob = "shaders/**/*.comp", accel = "vulkan1.2" }, +] + +[targets.spirv-module-consumer] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/spirv-module-consumer/shaders/a/scale.comp b/tests/spirv-module-consumer/shaders/a/scale.comp new file mode 100644 index 0000000..360eb26 --- /dev/null +++ b/tests/spirv-module-consumer/shaders/a/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-module-consumer/shaders/b/scale.comp b/tests/spirv-module-consumer/shaders/b/scale.comp new file mode 100644 index 0000000..360eb26 --- /dev/null +++ b/tests/spirv-module-consumer/shaders/b/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-module-consumer/src/main.cpp b/tests/spirv-module-consumer/src/main.cpp new file mode 100644 index 0000000..4897a5c --- /dev/null +++ b/tests/spirv-module-consumer/src/main.cpp @@ -0,0 +1,44 @@ +// Nothing here names a generated file. The shaders arrive through one import, +// and the two that share a stem are told apart by the directory they came from. +import std; +import spirv_module_consumer.shaders; + +namespace shaders = spirv_module_consumer::shaders; + +namespace { + +// The one property that identifies a SPIR-V module without executing it. +constexpr std::uint32_t kSpirvMagic = 0x07230203u; + +bool check(std::string_view label, shaders::payload p) { + const bool ok = p.size_bytes >= 4 && p.code[0] == kSpirvMagic; + std::cout << std::format("{}: magic={:08x} bytes={} {}\n", label, + p.size_bytes >= 4 ? p.code[0] : 0u, p.size_bytes, + ok ? "ok" : "BAD"); + return ok; +} + +} // namespace + +int main() { + bool ok = true; + ok &= check("a/scale.comp", shaders::a::scale_comp()); + ok &= check("b/scale.comp", shaders::b::scale_comp()); + + // The two shaders are the same source in two directories, so they must + // produce identical modules -- and they must be two distinct objects, not + // one symbol reached twice. Both are properties of the naming, and a build + // that collapsed them would still print two lines above. + const auto x = shaders::a::scale_comp(); + const auto y = shaders::b::scale_comp(); + if (x.code == y.code) { + std::cout << "BAD: both namespaces resolve to one array\n"; + ok = false; + } + if (x.size_bytes != y.size_bytes) { + std::cout << "BAD: identical sources produced different sizes\n"; + ok = false; + } + std::cout << (ok ? "all ok\n" : "FAILED\n"); + return ok ? 0 : 1; +} diff --git a/tools/embed.cppm b/tools/embed.cppm index 6a1dd82..eb562d8 100644 --- a/tools/embed.cppm +++ b/tools/embed.cppm @@ -31,6 +31,11 @@ export module mcpp.tools.embed; import std; import mcpp; +// The lib root, which carries `mcpp::plugins::surface` -- the declarations a +// consumer names. `group()` below hands its payloads to it, so a set of files +// embedded by this tool and a set of shaders compiled by `mcpp.rules.spirv` +// reach a consumer through the same shape. +import mcpp.plugins; // WHY NOTHING HERE USES `std::println`, AND WHY THAT IS NOT A STYLE CHOICE. @@ -215,4 +220,83 @@ inline bool files(std::span inputs, options opt = {}) { return true; } +// Several files, reached through ONE declaration a consumer imports. +// +// `files()` writes a header per input and leaves the consumer to include each +// by name. `group()` writes those same headers and then hands them to +// `mcpp::plugins::surface`, so the consumer writes one `import` and names no +// generated file -- the same surface `mcpp.rules.spirv` produces, from the same +// generator, because a payload that was already on disk and one a compiler +// produced are the same thing to whoever consumes it. +// +// The group's own name is required rather than derived. A rule knows what its +// payloads are for and can name the module `.shaders`; a tool called +// on an arbitrary set of files does not, and a derived name would be a guess +// that two calls in one build program could collide on. +inline bool group(std::span inputs, + const std::string& module_name, + mcpp::plugins::surface::kind surface + = mcpp::plugins::surface::default_surface(), + options opt = {}) { + if (inputs.empty()) return true; + if (module_name.empty()) { + std::cerr << "mcpp.tools.embed: group() needs a module name; it is what a " + "consumer imports and the namespace the declarations sit in\n"; + return false; + } + if (!files(inputs, opt)) return false; + + const auto dir = opt.out_dir.empty() ? default_dir() : opt.out_dir; + std::vector items; + for (auto const& one : inputs) { + const auto absolute = std::filesystem::path(one).is_absolute() + ? std::filesystem::path(one) + : std::filesystem::path(mcpp::manifest_dir()) / one; + const auto id = identifier_for(absolute, opt); + // `files()` wrote `.h` beside its siblings, so the include is the + // bare name: this tool's generated tree is flat, unlike a rule's, which + // mirrors the source tree it globbed. + // + // The symbol is QUALIFIED by `options::name_space`, because that is + // where `file()` put the array. Passing the bare name compiles for a + // caller that left the option empty and fails for one that did not, + // which is the shape of a defect that only the second test finds. + const auto sym = opt.name_space.empty() ? id : opt.name_space + "::" + id; + items.push_back({ .identifier = id, + .name_space = {}, + .data_header = id + ".h", + .data_symbol = sym, + // `_size` IS A COUNT OF ELEMENTS, AND THE SURFACE + // REPORTS BYTES. For `element::byte_` the two are the + // same number and the distinction is invisible; for + // `word32` it is four times out. + // + // `sizeof` is not the answer either: `null_terminate` + // appends a zero byte that `_size` deliberately does + // not count, so `sizeof` is one too many. Measured on + // a group of null-terminated text payloads, which + // reported one extra byte and failed the fixture's + // comparison on a trailing NUL. + .data_size_expr = opt.elem == element::word32 + ? sym + "_size * 4" + : sym + "_size" }); + } + + mcpp::plugins::surface::options so; + so.surface = surface; + so.elem = opt.elem == element::word32 + ? mcpp::plugins::surface::element::word32 + : mcpp::plugins::surface::element::byte_; + so.module_name = module_name; + so.out_dir = dir; + so.produced_by = "mcpp.tools.embed"; + + const auto out = mcpp::plugins::surface::emit(items, so); + if (!out) return false; + mcpp::generated(out->interface_file.c_str()); + mcpp::generated(out->impl_file.c_str()); + if (!out->include_dir.empty()) mcpp::include_dir(out->include_dir.c_str()); + return true; +} + } // namespace mcpp::tools::embed From 5f8cb460555af4b9cb0cd08f4373332b76a06f49 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Mon, 7 Sep 2026 22:45:23 +0800 Subject: [PATCH 02/20] Pin the engine that escapes for cmd.exe, and restore the floors it makes expressible `MCPP_VERSION` moves to 2026.9.7.1, the release carrying three things this package needs: `.slang` in the device-source extension table, without which `rules.slang` cannot route a file at all; `MCPP_LANGUAGE_MODULES`, which the surface's default reads; and `mcpp.platform.shell`, which escapes a provisioning argument for both cmd.exe and the child's MSVCRT parser. The third restores a constraint. `xim:shaderc` was an exact version on macOS and Windows for one release because the `>` in `>=2026.3` was read by cmd.exe as a redirection and arrived as a packaging failure; the comment in mcpp.toml said both would revert once a released engine escaped the argument. They do here. `xim:slang` becomes a floor for the same reason, and its comment records why the obvious middle ground was rejected: `>=` on the two platforms that tolerated it and an exact version on the third would resolve different payloads for one project depending on the machine that built it, which is how a defect comes to exist on one operating system only. --- .github/workflows/ci.yml | 2 +- mcpp.toml | 50 +++++++++++++++++++--------------------- rules/slang.cppm | 16 ++++++++++++- 3 files changed, 40 insertions(+), 28 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 20777a8..f29eac1 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -9,7 +9,7 @@ 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.6.6 + MCPP_VERSION: 2026.9.7.1 # PINNED, AND WITHOUT IT THE CACHE BELOW CACHED NOTHING. # # A released mcpp is self-contained: with no `MCPP_HOME`, `mcpp self env` diff --git a/mcpp.toml b/mcpp.toml index 8331f9e..e3552bb 100644 --- a/mcpp.toml +++ b/mcpp.toml @@ -140,14 +140,15 @@ tools-embed = { sources = ["tools/embed.cppm"] } [target.'cfg(all(accelerator = "vulkan", linux))'.feature-xlings.rules-spirv] "xim:glslang" = ">=15.1.0" -# AN EXACT VERSION HERE, WHERE LINUX HAS A FLOOR, AND THE REASON IS THE -# ENGINE'S COMMAND LINE RATHER THAN THE PACKAGE. +# A FLOOR AGAIN, NOW THAT THE ENGINE ESCAPES THE ARGUMENT FOR cmd.exe. # -# mcpp hands its provisioning request to xlings as a JSON argument on a shell -# command line. On Windows that line is parsed by cmd.exe, whose quoting rules -# are not the ones the JSON is escaped for: the `\"` sequences leave cmd's -# quote state OFF by the time it reaches a `>` in a version constraint, so the -# `>` is read as a REDIRECTION. Measured on windows-2022 with mcpp 2026.9.6.6: +# These two were an exact version for one release, and the reason was the +# engine's command line rather than the package. mcpp hands its provisioning +# request to xlings as a JSON argument on a shell command line; on Windows that +# line is parsed by cmd.exe, whose quoting rules are not the ones the JSON is +# escaped for. The `\"` sequences left cmd's quote state OFF by the time it +# reached the `>` in a version constraint, so the `>` was read as a REDIRECTION. +# Measured on windows-2022 with mcpp 2026.9.6.6: # # Provisioning [xlings.workspace] entries declared by dependencies # (xim:shaderc@>=2026.3) @@ -158,16 +159,14 @@ tools-embed = { sources = ["tools/embed.cppm"] } # packaging failure. No declaration reachable on Windows had carried a `>` # before, which is why the whole `>=` shape had never been exercised there. # -# An exact version is a legitimate declaration and not a workaround in -# disguise: mcpp reads it as a CHOICE, so a project that wants a different one -# still wins and the override is reported. macOS takes the same value so the -# two platforms that use this compiler agree. Both revert to `>=2026.3` once a -# released engine escapes the argument for cmd. +# `mcpp.platform.shell` now escapes for both parsers, and MCPP_VERSION below +# pins the release that carries it, so the constraint these entries were always +# meant to state is expressible again. [target.'cfg(all(accelerator = "vulkan", macos))'.feature-xlings.rules-spirv] -"xim:shaderc" = "2026.3" +"xim:shaderc" = ">=2026.3" [target.'cfg(all(accelerator = "vulkan", windows))'.feature-xlings.rules-spirv] -"xim:shaderc" = "2026.3" +"xim:shaderc" = ">=2026.3" # ONE TABLE FOR ALL THREE PLATFORMS, WHICH `rules-spirv` COULD NOT HAVE. # @@ -175,19 +174,18 @@ tools-embed = { sources = ["tools/embed.cppm"] } # aarch64, so this rule does not change compilers by platform the way the GLSL # rule does, and one table says everything. # -# AN EXACT VERSION RATHER THAN A FLOOR, AND NOT FOR THE USUAL REASON. Slang's -# version is coupled to no driver, so the rule this package follows would make -# it a `>=`. It cannot be one on Windows: mcpp hands the provisioning request to -# xlings as JSON on a cmd.exe command line, where the `>` is read as a -# redirection (see the note above `rules-spirv`'s Windows entry). -# -# Writing `>=` on the two platforms that tolerate it and an exact version on the -# third would resolve DIFFERENT payloads for one project depending on which -# machine built it, which is how a defect comes to exist on one operating system -# only. A uniform exact version is the weaker constraint honestly applied -# everywhere; a project that wants a newer Slang names it and wins. +# A FLOOR, WHICH IS WHAT THIS PACKAGE'S RULE FOR THE SHAPE ASKS FOR: Slang's +# version is coupled to no driver, so newer is simply newer. +# +# It is written once rather than per platform because the engine escapes the +# constraint for cmd.exe (see the note above), so the `>` survives on all three. +# An earlier revision of this file pinned it exactly on every platform to avoid +# that defect, and rejected the obvious alternative of `>=` on the two platforms +# that tolerated it: a floor holding on two of three would resolve DIFFERENT +# payloads for one project depending on the machine that built it, which is how +# a defect comes to exist on one operating system only. [target.'cfg(accelerator = "vulkan")'.feature-xlings.rules-slang] -"xim:slang" = "2026.14.1" +"xim:slang" = ">=2026.14.1" # 8.5.0 is a real floor rather than a preference: the mixed-mode object -- the # one carrying both the device binary and a host-callable launcher, which is diff --git a/rules/slang.cppm b/rules/slang.cppm index 0e8d0b6..7c9bad6 100644 --- a/rules/slang.cppm +++ b/rules/slang.cppm @@ -378,7 +378,21 @@ inline bool compile(std::span shaders, options opt = {}) { std::filesystem::create_directories(dir, ec); const auto base = (dir / p.stem().string()).string(); const auto header = base + ".h"; - const auto inc = base + ".inc"; + // SLANGC APPENDS THE EMBEDDING LANGUAGE'S EXTENSION TO `-o`. + // + // `-o scale.inc` writes `scale.inc.h`, because `-source-embed-language` + // defaults to C/C++ and the driver adds that language's suffix unless + // the name already ends in it. An action whose declared output is + // `scale.inc` therefore names a file the command never writes, and the + // failure lands two edges away: the generated implementation includes + // the wrapper, the wrapper includes a file that is not there, and the + // message is `fatal error: scale.inc: No such file or directory` with + // nothing pointing at the flag that caused it. + // + // Measured against slangc 2026.14.1. Naming the output `.h` outright + // makes the file the compiler writes and the file this rule declares + // the same one, without depending on that appending rule at all. + const auto inc = base + "_embed.h"; const auto input = p.is_absolute() ? src : root + "/" + src; if (!write_header(header, inc)) return false; From 9b897572264cd0258f33fae04f6ddade87e2235d Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Mon, 7 Sep 2026 23:41:41 +0800 Subject: [PATCH 03/20] Declare device_extensions and rule_module per rule feature A consumer now writes one dependency edge and no build program. The two keys are what the engine reads: which device sources a rule compiles, and which module a build program imports to reach it. The engine holds neither this package's name nor any feature spelling, so a third-party rule for the same extension works the same way. tests/spirv-zero-config is the whole path: an edge naming a feature, no host-module, no build.mcpp, and a shader reached through a generated module. --- mcpp.toml | 74 +++++++++++++++++++--- tests/spirv-zero-config/mcpp.toml | 43 +++++++++++++ tests/spirv-zero-config/shaders/scale.comp | 15 +++++ tests/spirv-zero-config/src/main.cpp | 13 ++++ 4 files changed, 137 insertions(+), 8 deletions(-) create mode 100644 tests/spirv-zero-config/mcpp.toml create mode 100644 tests/spirv-zero-config/shaders/scale.comp create mode 100644 tests/spirv-zero-config/src/main.cpp diff --git a/mcpp.toml b/mcpp.toml index e3552bb..ec50185 100644 --- a/mcpp.toml +++ b/mcpp.toml @@ -19,15 +19,73 @@ import_std = true [build] sources = ["src/plugins.cppm"] +# EACH RULE STATES WHAT IT COMPILES AND HOW TO REACH IT (mcpp 2026.9.7.1+). +# +# `device_extensions` says which device sources this rule compiles; +# `rule_module` says which module a consumer's build program imports to reach +# it. Two things follow, and neither puts this package's name inside mcpp: +# +# 1. A consumer that activates the feature gets those extensions classified +# as device sources, so a NEW device language costs no engine change. +# Slang measured the alternative: adding `.slang` to the engine's built-in +# table cost an mcpp release and a version bump in this file's CI before +# the rule could route one file. +# 2. `host-module = true` is implied, because a feature naming a rule module +# has already said that is the only way to use it. A consumer writes +# `features = ["rules-spirv"]` and nothing else. +# +# The feature is still requested BY NAME. An earlier revision activated it from +# the extensions a project's sources happened to carry; that was withdrawn +# because two packages may claim one extension -- a third-party CUDA rule is a +# thing someone will write -- and because a derived feature set is information +# a manifest no longer states. +# +# A consumer that outgrows the generated build program writes its own +# `build.mcpp`, which is the same program with edits. Each layer overrides the +# one above it; none is a different mechanism. [features] -default = [] -rules-ascendc = { sources = ["rules/ascendc.cppm"] } -rules-cuda = { sources = ["rules/cuda.cppm"] } -rules-hip = { sources = ["rules/hip.cppm"] } -rules-slang = { sources = ["rules/slang.cppm"] } -rules-spirv = { sources = ["rules/spirv.cppm"] } -rules-sycl = { sources = ["rules/sycl.cppm"] } -tools-embed = { sources = ["tools/embed.cppm"] } +default = [] + +[features.rules-ascendc] +sources = ["rules/ascendc.cppm"] +rule_module = "mcpp.rules.ascendc" +device_extensions = [".asc", ".cce"] + +[features.rules-cuda] +sources = ["rules/cuda.cppm"] +rule_module = "mcpp.rules.cuda" +device_extensions = [".cu"] + +[features.rules-hip] +sources = ["rules/hip.cppm"] +rule_module = "mcpp.rules.hip" +device_extensions = [".hip"] + +[features.rules-slang] +sources = ["rules/slang.cppm"] +rule_module = "mcpp.rules.slang" +device_extensions = [".slang"] + +# `.glsl` and `.hlsl` carry no stage and are claimed on purpose: the rule +# refuses them by name and says which extensions do carry one, which is a better +# message than the engine's "no rule compiles it". +[features.rules-spirv] +sources = ["rules/spirv.cppm"] +rule_module = "mcpp.rules.spirv" +device_extensions = [".comp", ".vert", ".frag", ".geom", ".tesc", ".tese", + ".mesh", ".task", ".rgen", ".rint", ".rahit", ".rchit", + ".rmiss", ".rcall", ".glsl", ".hlsl"] + +[features.rules-sycl] +sources = ["rules/sycl.cppm"] +rule_module = "mcpp.rules.sycl" +device_extensions = [".sycl"] + +# NOT a device rule: it embeds files that are already on disk, so it claims no +# extension and names no rule module. A consumer calls it from its own +# `build.mcpp`. +[features.tools-embed] +sources = ["tools/embed.cppm"] # ── The environment each rule needs (mcpp 2026.9.6.6+) ────────────────────── # diff --git a/tests/spirv-zero-config/mcpp.toml b/tests/spirv-zero-config/mcpp.toml new file mode 100644 index 0000000..d6f93ca --- /dev/null +++ b/tests/spirv-zero-config/mcpp.toml @@ -0,0 +1,43 @@ +# The consumer that has NO build.mcpp. +# +# `spirv-consumer` and `spirv-module-consumer` both write one, so neither can +# tell whether mcpp would have written it for them. This fixture is the whole +# path: a dependency edge naming a feature, and nothing else. +# +# What it asserts: +# +# - `host-module = true` is not written and is not needed. `rules-spirv` +# declares a `rule_module`, which is the statement that the only way to use +# the feature is as a host module. +# - No `build.mcpp` exists in this directory, and the shaders are still +# compiled: mcpp writes the program the feature describes into the build +# directory. +# - The surface defaults to the module form, because `[language] modules` is +# true and mcpp reports that as `MCPP_LANGUAGE_MODULES`. Nothing here asks +# for it. +[package] +name = "spirv-zero-config" +namespace = "example" +version = "0.1.0" +description = "A GLSL compute shader compiled with no build program of its own" + +[language] +standard = "c++23" +modules = true +import_std = true + +# The whole declaration. No `features` beyond the one that names the rule, no +# `host-module`, no build program. +[build-dependencies.mcpp] +plugins = { path = "../..", features = ["rules-spirv"] } + +[build] +accel = "vulkan1.2" +sources = [ + "src/*.cpp", + { glob = "shaders/*.comp", accel = "vulkan1.2" }, +] + +[targets.spirv-zero-config] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/spirv-zero-config/shaders/scale.comp b/tests/spirv-zero-config/shaders/scale.comp new file mode 100644 index 0000000..360eb26 --- /dev/null +++ b/tests/spirv-zero-config/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-zero-config/src/main.cpp b/tests/spirv-zero-config/src/main.cpp new file mode 100644 index 0000000..c702e65 --- /dev/null +++ b/tests/spirv-zero-config/src/main.cpp @@ -0,0 +1,13 @@ +// No build program in this project, and no generated file named here either. +import std; +import spirv_zero_config.shaders; + +int main() { + const auto s = spirv_zero_config::shaders::scale_comp(); + const bool ok = s.size_bytes >= 4 && s.code[0] == 0x07230203u; + std::cout << std::format("magic={:08x} bytes={} {}\n", + s.size_bytes >= 4 ? s.code[0] : 0u, s.size_bytes, + ok ? "ok" : "BAD"); + std::cout << (ok ? "all ok\n" : "FAILED\n"); + return ok ? 0 : 1; +} From 4444d5e22691f8d6b9484fc493703cdd0456674f Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Mon, 7 Sep 2026 23:48:59 +0800 Subject: [PATCH 04/20] Three storages behind one surface, each with a fixture that runs `storage` decides where a payload's bytes live; the surface decides how a consumer names them. The declarations are identical under all three, so a project changes the first and no consumer changes -- which is what the three new fixtures assert by having the same `src/main.cpp` shape. header a C array in generated source, compiled in (default) object a section, through `.incbin` in a generated .S sidecar a file beside the artifact, read at run time WHICH ONE IS A MEASUREMENT. With GCC 16.1 on 100 payloads of 16 KB each -- an ordinary compute shader's size -- the header route takes 1.10s against the object route's 2.17s, because at that size neither has a measurable marginal cost and the total is decided by how many processes start. The crossover is the TOTAL embedded byte count rather than the payload count: below about 1 MB the header route wins, above about 4 MB it loses by an order of magnitude (2.31s against 0.116s). So `header` stays the default and `object` is what a project reaches for when it has more payload than that. Object storage removes the one place the two shader compilers differed. Under it neither `-mfmt=c` nor `-x --vn` is passed, so glslang and glslc write the same thing: a bare `.spv`. The wrapper header exists only for the storage that needs a translation unit. The assembly is per object format, and the differences are not cosmetic: ELF takes `.section .rodata`, Mach-O `.section __TEXT,__const` with an underscore PREFIXED to every symbol, COFF `.section .rdata,"dr"`. Mach-O's underscore is the one that fails in the quiet direction -- a label without it defines a symbol the C++ side never resolves, and the link error names the accessor rather than the missing prefix. `.balign 4` is written because `VkShaderModuleCreateInfo:: pCode` requires it and a section directive alone does not promise it. MSVC has no GAS and mcpp refuses `.S` under it, so object storage degrades to header storage there with one warning. The surface is unchanged, so a consumer compiled either way is the same source. Sidecar states its cost instead of hiding it: the accessor opens a path relative to the working directory. `tests/spirv-sidecar` asserts both halves -- found from the package root, and reported missing when run from `/tmp`, which is the leg that would otherwise pass against an implementation that had embedded the payload after all. --- .github/workflows/ci.yml | 79 +++++ README.md | 37 +++ rules/spirv.cppm | 75 ++++- src/plugins.cppm | 272 ++++++++++++++++-- tests/spirv-object-storage/build.mcpp | 23 ++ tests/spirv-object-storage/mcpp.toml | 42 +++ tests/spirv-object-storage/shaders/scale.comp | 15 + tests/spirv-object-storage/src/main.cpp | 22 ++ tests/spirv-sidecar/build.mcpp | 24 ++ tests/spirv-sidecar/mcpp.toml | 42 +++ tests/spirv-sidecar/shaders/scale.comp | 15 + tests/spirv-sidecar/src/main.cpp | 21 ++ 12 files changed, 628 insertions(+), 39 deletions(-) create mode 100644 tests/spirv-object-storage/build.mcpp create mode 100644 tests/spirv-object-storage/mcpp.toml create mode 100644 tests/spirv-object-storage/shaders/scale.comp create mode 100644 tests/spirv-object-storage/src/main.cpp create mode 100644 tests/spirv-sidecar/build.mcpp create mode 100644 tests/spirv-sidecar/mcpp.toml create mode 100644 tests/spirv-sidecar/shaders/scale.comp create mode 100644 tests/spirv-sidecar/src/main.cpp diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index f29eac1..0d2077e 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -198,6 +198,85 @@ jobs: rm -f dup.log restored.log echo "ok: and it builds again with one shader per stem" + # NO BUILD PROGRAM AT ALL, WHICH IS WHAT THE TWO KEYS BUY. + # + # `spirv-consumer` and `spirv-module-consumer` both write a `build.mcpp`, + # so neither can tell whether mcpp would have written it. This fixture has + # a dependency edge naming a feature and nothing else: no `host-module`, + # no build program. The step asserts both absences, because a fixture that + # gained one by accident would still pass every other check here. + - name: rules-spirv with no build program + working-directory: tests/spirv-zero-config + run: | + test ! -f build.mcpp || { echo "FAIL: the fixture has its own build program"; exit 1; } + grep -q 'host-module' mcpp.toml && { echo "FAIL: the fixture declares host-module"; exit 1; } + "$MCPP" build | tee build.log + grep -q 'Rules mcpp.rules.spirv' build.log \ + || { echo "FAIL: the build did not report which rule it ran"; exit 1; } + "$MCPP" run | tee run.log + grep -q '^magic=07230203' run.log + grep -q '^all ok' run.log + # And the program mcpp wrote is where it says it is, and is the one a + # project would take over. + test -f target/.build-mcpp/build.mcpp + grep -q '^import mcpp.rules.spirv;' target/.build-mcpp/build.mcpp + echo "ok: one edge, no program, shaders compiled" + + # WHERE THE BYTES LIVE, TWICE, AGAINST ONE UNCHANGED CONSUMER. + # + # Both fixtures' `src/main.cpp` has the same shape as the header-stored + # one: one import, one call, a `payload` struct. That is the claim -- the + # storage is not something a consumer can see -- and it is why these are + # separate fixtures rather than options on an existing one. + # + # `object` is what a project reaches for above roughly 1 MB of total + # payload; below that the header route is faster, measured at 1.10s + # against 2.17s for 100 shaders of 16 KB. + - name: rules-spirv with object storage + working-directory: tests/spirv-object-storage + run: | + "$MCPP" build + "$MCPP" run | tee run.log + grep -q '^magic=07230203' run.log + grep -q '^all ok' run.log + # The bytes are a section, so the compiler never saw a C array: the + # payload is a bare `.spv` and an `.S` names it with `.incbin`. + test -f target/.build-mcpp/out/spirv/scale_comp.spv + test -f target/.build-mcpp/out/spirv/spirv_object_storage.shaders.payload.S + grep -q '\.incbin' target/.build-mcpp/out/spirv/spirv_object_storage.shaders.payload.S + if ls target/.build-mcpp/out/spirv/*.inc >/dev/null 2>&1; then + echo "FAIL: object storage still produced a C initialiser list" + exit 1 + fi + # `pCode` wants four-byte alignment and a section directive alone does + # not promise it. + grep -q '\.balign 4' target/.build-mcpp/out/spirv/spirv_object_storage.shaders.payload.S + echo "ok: the payload is a section, not generated source" + + - name: rules-spirv with sidecar storage + working-directory: tests/spirv-sidecar + run: | + "$MCPP" build + "$MCPP" run | tee run.log + grep -q '^magic=07230203' run.log + grep -q '^all ok' run.log + test -f target/.build-mcpp/out/spirv/scale_comp.spv + # THE REVERSE LEG, AND IT IS THE POINT OF THIS STORAGE. The accessor + # opens a path relative to the working directory, so the program finds + # its shaders from the package root and must not from anywhere else. + # Without this half the fixture would pass against an implementation + # that had quietly embedded the payload after all. + bin=$(ls target/*/*/bin/spirv-sidecar) + bin=$(cd "$(dirname "$bin")" && pwd)/$(basename "$bin") + if (cd /tmp && "$bin" > away.log 2>&1); then + echo "FAIL: the sidecar program succeeded from the wrong directory" + cat /tmp/away.log + exit 1 + fi + grep -q 'sidecar payload was not found' /tmp/away.log \ + || { echo "FAIL: it failed for some other reason"; cat /tmp/away.log; exit 1; } + echo "ok: found from the package root, reported missing from elsewhere" + # SLANG: A DIFFERENT LANGUAGE, THE SAME SURFACE. # # The point of this step is not that Slang compiles -- it is that a diff --git a/README.md b/README.md index 44dab72..6e78f68 100644 --- a/README.md +++ b/README.md @@ -168,6 +168,43 @@ std-free equivalent, and that cost is fixed rather than proportional to the payload -- it is ``'s templates, present whether the payload is 16 KB or 16 MB. A consumer that wants a `std::span` constructs one from the two members. +**Where the bytes live is a second, independent choice.** The surface decides how +a consumer names a payload; `storage` decides where it sits. The declarations are +identical under all three, so a project changes this and no consumer changes. + +| storage | the payload is | reach for it when | +|---|---|---| +| `header` (default) | a C array in generated source, compiled in | almost always | +| `object` | a section, through `.incbin` in a generated `.S` | total payload is large | +| `sidecar` | a file beside the artifact, read at run time | hot reload, or a payload too large to link | + +**Which one is a measurement, not a preference.** With GCC 16.1 on 100 payloads +of 16 KB each -- the size of an ordinary compute shader: + +``` +header route compile 0.64s + link 0.44s = 1.10s +object route convert 1.28s + compile 0.44s + link 0.46s = 2.17s +``` + +The header route is faster, because at that size neither route has a measurable +marginal cost and the total is decided by how many processes start; one compiler +invocation absorbs many headers. The crossover is the TOTAL embedded byte count +rather than the payload count: below about 1 MB the header route wins, and above +about 4 MB the compiler's slightly superlinear curve loses by an order of +magnitude (2.31s against 0.116s). Source expansion is a constant 2.75x. + +**`object` needs a GAS assembler.** Every gcc and clang toolchain has one on all +three platforms; MSVC does not, and mcpp refuses `.S` under it, so the emitter +falls back to `header` there and says so once. The surface does not change, so a +consumer compiled either way is the same source. + +**`sidecar` states its cost rather than hiding it.** The accessor opens a path +relative to the working directory, so the program finds its payloads when run +from the package root and does not when run from elsewhere -- which is why it is +not the default, and why `mcpp pack` of such a program has something further to +collect. `tests/spirv-sidecar` asserts both halves: found from the root, and +reported missing from `/tmp`. + **The default follows the project.** `[language] modules = true` gives the module surface, `false` gives a header with the same declarations. mcpp reports the setting as `MCPP_LANGUAGE_MODULES`; an engine that does not report it leaves diff --git a/rules/spirv.cppm b/rules/spirv.cppm index 3823726..2bdeda5 100644 --- a/rules/spirv.cppm +++ b/rules/spirv.cppm @@ -113,6 +113,17 @@ struct options { // every consumer of this package had before the surface existed. mcpp::plugins::surface::kind surface = mcpp::plugins::surface::default_surface(); + // Where the compiled SPIR-V lives. `header` compiles it in as generated + // source, `object` as a section reached through `.incbin`, `sidecar` as a + // file beside the artifact. See `mcpp::plugins::surface::storage` for the + // measurement that makes `header` the default. + // + // It changes what this rule asks the compiler for. Under `header` the + // compiler is told to emit a C declaration (`-mfmt=c`, `-x --vn`); under + // the other two it emits a bare `.spv`, which is both simpler and the one + // shape both compilers agree on. + mcpp::plugins::surface::storage storage = mcpp::plugins::surface::storage::header; + // The module a consumer imports, and the namespace the declarations sit in: // `myapp.shaders` gives `myapp::shaders::blur_comp()`. Empty derives it from // the package directory, so a project that states nothing still gets a name @@ -725,15 +736,32 @@ inline bool compile(std::span shaders, options opt = {}) { std::filesystem::create_directories(dir, ec); const auto base = (dir / (p.stem().string() + "_" + std::string(stage))).string(); const auto header = base + ".h"; + // Declared here rather than beside the action below, because the item + // recorded a few lines down names it: the surface needs the payload's + // path to write `.incbin`, and it writes that before any action runs. + const std::string spv = base + ".spv"; // As an `#include` writes it: relative to `gen`, which is the directory // this rule puts on the include path. std::string headerRel; for (auto const& seg : ns) headerRel += seg + "/"; headerRel += p.stem().string() + "_" + std::string(stage) + ".h"; - items.push_back({ .identifier = p.stem().string() + "_" + std::string(stage), - .name_space = ns, - .data_header = headerRel, - .data_symbol = sym }); + // WHERE A SIDECAR IS FOUND AT RUN TIME, AND WHAT THAT COSTS. + // + // The generated accessor opens this path relative to the WORKING + // DIRECTORY, so it is written relative to the package root -- which is + // where `mcpp run` starts the program, and where a project doing shader + // hot-reload runs it from. A program started from anywhere else finds + // nothing, and that is the property that keeps this storage off the + // default rather than a defect in it: the payload is not in the + // artifact, so something outside the artifact has to be true. + const auto sidecarName = + std::filesystem::path(spv).lexically_relative(root).generic_string(); + items.push_back({ .identifier = p.stem().string() + "_" + std::string(stage), + .name_space = ns, + .data_header = headerRel, + .data_symbol = sym, + .payload_path = spv, + .sidecar_name = sidecarName }); const auto input = std::filesystem::path(src).is_absolute() ? src : root + "/" + src; @@ -746,10 +774,15 @@ inline bool compile(std::span shaders, options opt = {}) { // rule writes `.h` around it. glslang used to write the header // itself, which made the two routes' headers differ in whether they // could be included first -- see `write_header`. + const bool embedAsSource = + opt.storage == mcpp::plugins::surface::storage::header; const std::string inc = base + ".inc"; - const std::string output = inc; + const std::string output = embedAsSource ? inc : spv; - if (!write_header(header, inc, sym, cc.kind)) return false; + // The wrapper header exists only to make the compiler's C output a + // translation unit. The other two storages never read a header, so + // writing one would leave a file nothing includes. + if (embedAsSource && !write_header(header, inc, sym, cc.kind)) return false; mcpp::action a; a.id = id.c_str(); @@ -777,16 +810,22 @@ inline bool compile(std::span shaders, options opt = {}) { 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()); - if (cc.kind == flavour::glslc) { - // `-mfmt=c` is the initialiser list; the declaration around it was - // written above. - a.arg("-mfmt=c"); - } else { - // `-x --vn` is what makes glslang's 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()); + if (embedAsSource) { + if (cc.kind == flavour::glslc) { + // `-mfmt=c` is the initialiser list; the declaration around it + // was written above. + a.arg("-mfmt=c"); + } else { + // `-x --vn` is what makes glslang's 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()); + } } + // Under `object` and `sidecar` neither flag is passed, so both + // compilers write the same thing: a bare SPIR-V module. The one place + // the two flavours differed disappears with the storage that needed it. a.arg("-o"); a.arg(output.c_str()); a.arg(input.c_str()); a.input(input.c_str()); @@ -802,6 +841,7 @@ inline bool compile(std::span shaders, options opt = {}) { // ── What a consumer names ──────────────────────────────────────────────── mcpp::plugins::surface::options so; so.surface = opt.surface; + so.store = opt.storage; so.elem = mcpp::plugins::surface::element::word32; so.module_name = moduleName; so.out_dir = gen; @@ -816,6 +856,11 @@ inline bool compile(std::span shaders, options opt = {}) { // declared about it. mcpp::generated(out->interface_file.c_str()); mcpp::generated(out->impl_file.c_str()); + // The `.S` under object storage. It is an ordinary source: mcpp assembles + // it, and `.incbin` reads the payload the action above produced, which by + // then exists because a `role = "source"` action is ordered before this + // package's compiles. + if (!out->assembly_file.empty()) mcpp::generated(out->assembly_file.c_str()); if (!out->include_dir.empty()) mcpp::include_dir(out->include_dir.c_str()); // Which collection produced these actions. The version was `0.1.1` while diff --git a/src/plugins.cppm b/src/plugins.cppm index 7ca8f0f..3bc2bf6 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -119,6 +119,45 @@ enum class kind { module_, c_header }; // call site, which is undefined behaviour on an under-aligned byte array. enum class element { byte_, word32 }; +// WHERE THE BYTES LIVE. Orthogonal to `kind`, which decides how a consumer +// NAMES them: the declarations are identical under all three, so a project +// changes this and no consumer changes. +// +// WHICH ONE TO USE IS A MEASUREMENT, NOT A PREFERENCE. With GCC 16.1 on this +// project's own fixtures, 100 payloads of 16 KB each: +// +// header route compile 0.64s + link 0.44s = 1.10s +// object route convert 1.28s + compile 0.44s + link 0.46s = 2.17s +// +// The header route is faster, because at a real shader's size neither route has +// a measurable marginal cost and the total is decided by how many processes +// start -- and one compiler invocation absorbs many headers. The crossover is +// the TOTAL embedded byte count, not the payload count: below about 1 MB the +// header route wins, and above about 4 MB the compiler's slightly superlinear +// curve loses by an order of magnitude (2.31s against 0.116s at 4 MB). Source +// expansion is a constant 2.75x, which is the second half of the same reason. +// +// So `header` is the default and `object` is what a project reaches for when it +// has more payload than that, not what it reaches for because objects sound +// tidier. +enum class storage { + // The payload is a C array in generated source, compiled into the program. + // No assembler involved, so it works on every toolchain including MSVC. + header, + // The payload is a section in an object, reached through `.incbin` in a + // generated `.S`. The bytes never pass through the C++ compiler. + // + // Requires a GAS-capable assembler. Every gcc and clang toolchain has one + // on all three platforms; MSVC does not, and mcpp refuses `.S` under it, so + // the emitter falls back to `header` there and says so once. + object, + // The payload is written beside the artifact and read at run time. The + // program's correctness then depends on its working directory, which is the + // reason this is not the default -- but it is what shader hot-reload needs, + // and what a payload too large to link needs. + sidecar, +}; + // One payload in a group. struct item { // The C++ identifier, already sanitised by the caller: `blur_comp`. @@ -136,7 +175,17 @@ struct item { // could not tell them apart. std::string data_header; // The array that header declares, qualified if it sits in a namespace. + // Unused under `storage::object`, where the linker symbol is derived from + // the accessor name instead, and under `storage::sidecar`. std::string data_symbol; + // The payload file itself, as an absolute path. Required by `object`, whose + // generated `.S` names it in `.incbin`, and by `sidecar`, which copies it. + // The file need not exist yet: `.incbin` resolves its argument at assembly + // time, which is after the action that writes it has run, and that is what + // lets the whole surface be written at plan time. + std::string payload_path; + // Where a sidecar payload is found at run time, relative to the artifact. + std::string sidecar_name; // The payload's size IN BYTES, as an expression valid where the generated // implementation writes it. Empty means `sizeof `. // @@ -152,6 +201,7 @@ struct item { struct options { kind surface = kind::module_; + storage store = storage::header; element elem = element::word32; // The module a consumer imports, e.g. `myapp.shaders`. The namespace is // this name with `.` replaced by `::`, which is why there is no second @@ -169,6 +219,13 @@ struct options { struct emitted { // The interface unit: a `.cppm` under `module_`, a `.h` under `c_header`. std::string interface_file; + // The generated `.S` under `storage::object`, empty otherwise. The caller + // adds it to the build the same way it adds the implementation. + std::string assembly_file; + // The storage actually used. Differs from what was asked for when the + // toolchain cannot assemble: MSVC has no GAS, so `object` degrades to + // `header` and this says so rather than leaving the caller to assume. + storage store = storage::header; // The translation unit defining the accessors. Always a plain `.cpp`. std::string impl_file; // Non-empty under `module_`: what the interface unit provides, which the @@ -304,6 +361,64 @@ inline std::string declarations(std::span items, const options& opt) return s; } +// ONE OBJECT FORMAT PER PLATFORM, AND THE DIFFERENCES ARE NOT COSMETIC. +// +// `.incbin` is the portable part: gas and clang's integrated assembler both +// accept it everywhere, and it resolves its argument at ASSEMBLY time, which is +// why this file can be written before the payload exists. What is not portable +// is the section directive and whether a C symbol carries a leading underscore. +// +// ELF `.section .rodata`, symbol as written +// Mach-O `.section __TEXT,__const`, symbol PREFIXED with `_` +// COFF `.section .rdata,"dr"`, symbol as written on x86_64 +// +// Mach-O's underscore is the one that fails quietly in the other direction: an +// assembly label without it defines a symbol the C++ side never resolves, and +// the link error names the accessor rather than the missing prefix. +struct asm_dialect { + std::string_view section; + std::string_view symbol_prefix; +}; + +inline asm_dialect dialect_for(std::string_view targetOs) { + if (targetOs == "macos" || targetOs == "macosx" || targetOs == "darwin") + return { ".section __TEXT,__const", "_" }; + if (targetOs == "windows") return { ".section .rdata,\"dr\"", "" }; + return { ".section .rodata", "" }; +} + +// `.balign 4` rather than nothing: `VkShaderModuleCreateInfo::pCode` requires +// four-byte alignment, and a section directive alone does not promise it. The +// header route gets alignment from the array's element type; this route has to +// ask for it. +inline std::string assembly_for(std::span items, const options& opt, + std::string_view targetOs) { + const auto d = dialect_for(targetOs); + std::string s = + "// Generated by mcpp.plugins.surface. Do not edit.\n" + "//\n" + "// The payloads, as sections rather than as C arrays. `.incbin` resolves\n" + "// its argument when this file is assembled, which is after the actions\n" + "// that write those files have run.\n"; + for (auto const& it : items) { + const auto base = accessor_base(opt, it); + s += std::format("\n {}\n" + " .globl {}{}_begin\n" + " .balign 4\n" + "{}{}_begin:\n" + " .incbin \"{}\"\n" + " .globl {}{}_end\n" + "{}{}_end:\n", + d.section, + d.symbol_prefix, base, + d.symbol_prefix, base, + it.payload_path, + d.symbol_prefix, base, + d.symbol_prefix, base); + } + return s; +} + // ---- the tool --------------------------------------------------------------- // Writes the interface and the implementation, and returns what the caller has @@ -324,10 +439,25 @@ inline std::optional emit(std::span items, const options& o const auto segs = split_module_name(opt.module_name); const auto by = opt.produced_by.empty() ? std::string("mcpp.plugins.surface") : opt.produced_by; - const auto dir = std::filesystem::path(opt.out_dir); - const auto body = declarations(items, opt); + // MSVC HAS NO GAS, AND mcpp REFUSES `.S` UNDER IT. + // + // `src/build/prepare.cppm` states that outright: "GAS assembly sources + // (.S/.s) are not supported by the MSVC toolchain". So object storage + // degrades to header storage there rather than producing a file the build + // will refuse. The SURFACE does not change -- the declarations are the same + // under both -- so a consumer compiled either way is the same source. emitted out; + out.store = opt.store; + if (out.store == storage::object + && std::string_view(mcpp::compiler()) == "msvc") { + out.store = storage::header; + mcpp::warning("mcpp.plugins.surface: object storage needs a GAS assembler and the " + "MSVC toolchain has none; the payload is compiled in as generated " + "source instead. The declarations a consumer sees are unchanged."); + } + const auto dir = std::filesystem::path(opt.out_dir); + const auto body = declarations(items, opt); // ---- interface ---- std::string iface; @@ -354,29 +484,113 @@ inline std::optional emit(std::span items, const options& o // ---- implementation ---- // - // The one translation unit that includes the data headers. It is a plain - // `.cpp` under both surfaces, because the accessors have C language linkage - // and so need no attachment to the module. - std::string impl; - impl += std::format("// Generated by mcpp.plugins.surface for {}. Do not edit.\n" - "//\n" - "// The only translation unit that includes the generated data headers.\n" - "// Each declares a `static` array, so this is also the only copy of the\n" - "// bytes in the program.\n", by); - for (auto const& it : items) impl += std::format("#include \"{}\"\n", it.data_header); - impl += "\n"; - impl += element_width_assertion(opt.elem); - impl += "\n"; + // The one translation unit that defines the accessors. A plain `.cpp` under + // every surface and every storage, because the accessors have C language + // linkage and so need no attachment to the module. What differs between the + // three storages is only where it reads the bytes from. const char* elem = element_type(opt.elem); - for (auto const& it : items) { - const auto base = accessor_base(opt, it); - const auto size = it.data_size_expr.empty() - ? std::format("sizeof {}", it.data_symbol) - : it.data_size_expr; - impl += std::format( - "extern \"C\" const {0}* {1}_data() {{ return {2}; }}\n" - "extern \"C\" unsigned long {1}_size() {{ return {3}; }}\n", - elem, base, it.data_symbol, size); + std::string impl; + impl += std::format("// Generated by mcpp.plugins.surface for {}. Do not edit.\n", by); + + if (out.store == storage::header) { + impl += "//\n" + "// The only translation unit that includes the generated data headers.\n" + "// Each declares a `static` array, so this is also the only copy of the\n" + "// bytes in the program.\n"; + for (auto const& it : items) impl += std::format("#include \"{}\"\n", it.data_header); + impl += "\n"; + impl += element_width_assertion(opt.elem); + impl += "\n"; + for (auto const& it : items) { + const auto base = accessor_base(opt, it); + const auto size = it.data_size_expr.empty() + ? std::format("sizeof {}", it.data_symbol) + : it.data_size_expr; + impl += std::format( + "extern \"C\" const {0}* {1}_data() {{ return {2}; }}\n" + "extern \"C\" unsigned long {1}_size() {{ return {3}; }}\n", + elem, base, it.data_symbol, size); + } + } else if (out.store == storage::object) { + impl += "//\n" + "// The bytes are in a section written by the generated `.S`. This file\n" + "// only names its two boundary symbols, so nothing here parses a payload\n" + "// and the C++ compiler never sees one.\n" + "//\n" + "// The symbols are declared as arrays of the element type rather than as\n" + "// `char`: the assembly aligned the section to four bytes, and a\n" + "// declaration that said `char` would let a consumer reach it through an\n" + "// under-aligned pointer with nothing to notice.\n"; + impl += "\n"; + impl += element_width_assertion(opt.elem); + impl += "\nextern \"C\" {\n"; + for (auto const& it : items) { + const auto base = accessor_base(opt, it); + impl += std::format("extern const {0} {1}_begin[];\n" + "extern const {0} {1}_end[];\n", elem, base); + } + impl += "}\n\n"; + for (auto const& it : items) { + const auto base = accessor_base(opt, it); + impl += std::format( + "extern \"C\" const {0}* {1}_data() {{ return {1}_begin; }}\n" + "extern \"C\" unsigned long {1}_size() {{\n" + " return static_cast(({1}_end - {1}_begin) * sizeof({0}));\n" + "}}\n", elem, base); + } + } else { + impl += "//\n" + "// The payloads are files beside the artifact, read on first use. The\n" + "// path is resolved against the WORKING DIRECTORY, which is the property\n" + "// that makes this storage the one a project opts into rather than the\n" + "// default: a program started from elsewhere finds nothing.\n" + "//\n" + "// Read once and kept: an accessor that reloaded would hand two callers\n" + "// two different pointers to the same payload, and a device API given\n" + "// the second after the first was freed is a defect with no message.\n"; + impl += "#include \n#include \n#include \n\n"; + impl += element_width_assertion(opt.elem); + impl += std::format(R"IMPL( +namespace {{ + +struct blob {{ {0}* data = nullptr; unsigned long size = 0; bool tried = false; }}; + +blob& load(const char* path, blob& b) {{ + if (b.tried) return b; + b.tried = true; + std::FILE* f = std::fopen(path, "rb"); + if (!f) {{ + std::fprintf(stderr, "mcpp.plugins.surface: cannot open %s\n", path); + return b; + }} + std::fseek(f, 0, SEEK_END); + const long n = std::ftell(f); + std::fseek(f, 0, SEEK_SET); + if (n > 0) {{ + b.data = static_cast<{0}*>(std::malloc(static_cast(n))); + if (b.data && std::fread(b.data, 1, static_cast(n), f) + == static_cast(n)) + b.size = static_cast(n); + else {{ std::free(b.data); b.data = nullptr; }} + }} + std::fclose(f); + return b; +}} + +}} // namespace +)IMPL", elem); + impl += "\n"; + for (auto const& it : items) { + const auto base = accessor_base(opt, it); + impl += std::format( + "static blob {0}_blob;\n" + "extern \"C\" const {1}* {0}_data() {{\n" + " return load(\"{2}\", {0}_blob).data;\n" + "}}\n" + "extern \"C\" unsigned long {0}_size() {{\n" + " return load(\"{2}\", {0}_blob).size;\n" + "}}\n", base, elem, it.sidecar_name); + } } out.impl_file = (dir / (opt.module_name + ".impl.cpp")).string(); @@ -384,6 +598,16 @@ inline std::optional emit(std::span items, const options& o std::cerr << std::format("mcpp.plugins.surface: cannot write {}\n", out.impl_file); return std::nullopt; } + + if (out.store == storage::object) { + out.assembly_file = (dir / (opt.module_name + ".payload.S")).string(); + if (!write_if_different(out.assembly_file, + assembly_for(items, opt, mcpp::target_os()))) { + std::cerr << std::format("mcpp.plugins.surface: cannot write {}\n", + out.assembly_file); + return std::nullopt; + } + } return out; } diff --git a/tests/spirv-object-storage/build.mcpp b/tests/spirv-object-storage/build.mcpp new file mode 100644 index 0000000..61954c2 --- /dev/null +++ b/tests/spirv-object-storage/build.mcpp @@ -0,0 +1,23 @@ +import std; +import mcpp; +import mcpp.plugins; +import mcpp.rules.spirv; + +// STORAGE IS THE ONE THING THIS FIXTURE CHANGES. +// +// The consumer's source is byte-identical to `spirv-module-consumer`'s idea of +// how a shader is reached: one import, one call, a `payload` struct. What moves +// is where the bytes live -- a section written by a generated `.S` rather than a +// C array in generated source -- and the point of the fixture is that moving it +// changes nothing a consumer can see. +// +// A project reaches for this above roughly 1 MB of total payload; below that the +// header storage is faster, for the reason `mcpp::plugins::surface::storage` +// records with the numbers. +int main() { + mcpp::rerun_if_changed_glob("shaders/**/*.comp"); + mcpp::rules::spirv::options opt; + opt.surface = mcpp::plugins::surface::kind::module_; + opt.storage = mcpp::plugins::surface::storage::object; + return mcpp::rules::spirv::compile(opt) ? 0 : 1; +} diff --git a/tests/spirv-object-storage/mcpp.toml b/tests/spirv-object-storage/mcpp.toml new file mode 100644 index 0000000..4894d83 --- /dev/null +++ b/tests/spirv-object-storage/mcpp.toml @@ -0,0 +1,42 @@ +# 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-object-storage" +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 + +# `[build-dependencies]`, not `[dependencies]`: a rule package's library must +# never reach the target while its rule is still wanted, which is the case +# docs/05 section 2.6.1 exists for. `host-module = true` says which build-time +# product is wanted; the section says whether the package reaches the target. +[build-dependencies.mcpp] +plugins = { path = "../..", features = ["rules-spirv"], host-module = true } + +# NO [xlings.workspace]. The rule declares the payloads it drives, under the +# feature that selects it and the accelerator it serves, so this project names +# the rule and nothing else. Needs mcpp 2026.9.6.6; before it, a payload a +# dependency declared was installed and then answered as absent. + +[build] +accel = "vulkan1.2" +sources = [ + "src/*.cpp", + { glob = "shaders/*.comp", accel = "vulkan1.2" }, +] + +[targets.spirv-object-storage] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/spirv-object-storage/shaders/scale.comp b/tests/spirv-object-storage/shaders/scale.comp new file mode 100644 index 0000000..360eb26 --- /dev/null +++ b/tests/spirv-object-storage/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-object-storage/src/main.cpp b/tests/spirv-object-storage/src/main.cpp new file mode 100644 index 0000000..23eefd8 --- /dev/null +++ b/tests/spirv-object-storage/src/main.cpp @@ -0,0 +1,22 @@ +// Identical in shape to the header-storage consumer: the storage is not +// something a consumer can see. +import std; +import spirv_object_storage.shaders; + +int main() { + const auto s = spirv_object_storage::shaders::scale_comp(); + const bool ok = s.size_bytes >= 4 && s.code[0] == 0x07230203u; + std::cout << std::format("magic={:08x} bytes={} {}\n", + s.size_bytes >= 4 ? s.code[0] : 0u, s.size_bytes, + ok ? "ok" : "BAD"); + + // SPIR-V is a sequence of 32-bit words. Under object storage the size comes + // from subtracting two linker symbols rather than from `sizeof`, so a + // section the assembler padded would show up here and nowhere else. + if (s.size_bytes % 4 != 0) { + std::cout << "BAD: size is not a whole number of SPIR-V words\n"; + return 1; + } + std::cout << (ok ? "all ok\n" : "FAILED\n"); + return ok ? 0 : 1; +} diff --git a/tests/spirv-sidecar/build.mcpp b/tests/spirv-sidecar/build.mcpp new file mode 100644 index 0000000..892f0c1 --- /dev/null +++ b/tests/spirv-sidecar/build.mcpp @@ -0,0 +1,24 @@ +import std; +import mcpp; +import mcpp.plugins; +import mcpp.rules.spirv; + +// THE THIRD STORAGE: the payload stays a file and the program reads it. +// +// The consumer's source is again unchanged -- one import, one call, a `payload` +// struct -- which is the whole claim the surface makes. What changes is that the +// bytes are not in the artifact at all. +// +// This is what shader hot reload needs, and what a payload too large to link +// needs. Its cost is stated rather than hidden: the accessor opens a path +// relative to the WORKING DIRECTORY, so the program finds its shaders when run +// from the package root and does not when run from elsewhere. That is why it is +// not the default, and it is also why `mcpp pack` of such a program has +// something further to collect where a header- or object-stored one has nothing. +int main() { + mcpp::rerun_if_changed_glob("shaders/**/*.comp"); + mcpp::rules::spirv::options opt; + opt.surface = mcpp::plugins::surface::kind::module_; + opt.storage = mcpp::plugins::surface::storage::sidecar; + return mcpp::rules::spirv::compile(opt) ? 0 : 1; +} diff --git a/tests/spirv-sidecar/mcpp.toml b/tests/spirv-sidecar/mcpp.toml new file mode 100644 index 0000000..b90e426 --- /dev/null +++ b/tests/spirv-sidecar/mcpp.toml @@ -0,0 +1,42 @@ +# 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-sidecar" +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 + +# `[build-dependencies]`, not `[dependencies]`: a rule package's library must +# never reach the target while its rule is still wanted, which is the case +# docs/05 section 2.6.1 exists for. `host-module = true` says which build-time +# product is wanted; the section says whether the package reaches the target. +[build-dependencies.mcpp] +plugins = { path = "../..", features = ["rules-spirv"], host-module = true } + +# NO [xlings.workspace]. The rule declares the payloads it drives, under the +# feature that selects it and the accelerator it serves, so this project names +# the rule and nothing else. Needs mcpp 2026.9.6.6; before it, a payload a +# dependency declared was installed and then answered as absent. + +[build] +accel = "vulkan1.2" +sources = [ + "src/*.cpp", + { glob = "shaders/*.comp", accel = "vulkan1.2" }, +] + +[targets.spirv-sidecar] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/spirv-sidecar/shaders/scale.comp b/tests/spirv-sidecar/shaders/scale.comp new file mode 100644 index 0000000..360eb26 --- /dev/null +++ b/tests/spirv-sidecar/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-sidecar/src/main.cpp b/tests/spirv-sidecar/src/main.cpp new file mode 100644 index 0000000..20bcc2f --- /dev/null +++ b/tests/spirv-sidecar/src/main.cpp @@ -0,0 +1,21 @@ +// The same shape as the header- and object-stored consumers. Nothing here knows +// that the payload is a file. +import std; +import spirv_sidecar.shaders; + +int main() { + const auto s = spirv_sidecar::shaders::scale_comp(); + if (s.code == nullptr || s.size_bytes == 0) { + // The failure this storage can have and the other two cannot: the file + // was not where the program looked. Reported here rather than left to + // crash, because "started from the wrong directory" is the whole of its + // contract and a null pointer says nothing. + std::cout << "BAD: the sidecar payload was not found\n"; + return 1; + } + const bool ok = s.size_bytes >= 4 && s.code[0] == 0x07230203u; + std::cout << std::format("magic={:08x} bytes={} {}\n", s.code[0], s.size_bytes, + ok ? "ok" : "BAD"); + std::cout << (ok ? "all ok\n" : "FAILED\n"); + return ok ? 0 : 1; +} From 956a17f7f7d6c13d42089be6a02bab6dd92a12a6 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Mon, 7 Sep 2026 23:53:40 +0800 Subject: [PATCH 05/20] README: the file-name to call mapping, and the generator's reuse boundary Every name a consumer writes is derived, and the derivation was documented nowhere. The table states it end to end, with the three invariants and the defect each one exists to prevent -- including the byte-identical generated headers that GCC's `#pragma once` collapsed into one. The generator is not shader-specific and three members already share it. A rule package outside this collection reaches it by depending on `mcpp:plugins` and activating no feature; whether it should become a package of its own is left open rather than answered by silence. --- README.md | 53 +++++++++++++++++++++++++++++++++++++++++++++++++---- 1 file changed, 49 insertions(+), 4 deletions(-) diff --git a/README.md b/README.md index 6e78f68..75a024c 100644 --- a/README.md +++ b/README.md @@ -154,10 +154,42 @@ const auto s = myapp::shaders::blur_comp(); VkShaderModuleCreateInfo ci{ .codeSize = s.size_bytes, .pCode = s.code }; ``` -**The module name and the namespace are one identifier path.** `myapp.shaders` -gives `myapp::shaders`, and a payload's directory below the globbed tree adds a -segment: `shaders/post/tone.comp` is `myapp::shaders::post::tone_comp()`. The -name is derived from the package unless the project sets `options::module_name`. +### From a file name to a call + +Every name a consumer writes is derived, and derived one way, so nothing has to +be looked up: + +``` +base directory shaders/ derived: the shallowest directory + every payload shares +file shaders/post/tone.frag + └── the path below the base +module myapp.shaders package name + group +namespace myapp::shaders::post the module name segment by segment, + then the directory's segments +identifier tone_frag stem + stage, non-identifier + characters replaced by `_` +call myapp::shaders::post::tone_frag() +``` + +`myapp` comes from the package unless the project sets `options::module_name`; +the base directory is derived unless it sets `options::base_dir`. + +Three invariants, and each exists because its absence was a defect: + +- **The module name and the namespace are the same identifier path**, `.` for + `::`. A reader never has to learn which namespace a module opens. +- **The directory reaches the generated file's path and the linker symbol, not + only the namespace.** Two shaders sharing a stem in different directories + produced byte-identical generated headers, and GCC's `#pragma once` treats two + files with the same size and content as the same file -- so the second include + did nothing and both accessors returned the first array, while the program + printed the right magic number twice. +- **The stage is always part of the identifier**, so `blur.comp` and `blur.frag` + do not collide. Uniformly rather than only when needed: conditional naming is + worse than verbose naming. + + **The interface is a function, and it names no standard-library type.** A variable cannot keep one shape across the ways bytes can be stored, because @@ -211,6 +243,19 @@ the setting as `MCPP_LANGUAGE_MODULES`; an engine that does not report it leaves the header surface in place, so an older engine keeps the behaviour every consumer of this package had before the surface existed. +**The generator is not shader-specific, and three members already share it.** +`mcpp::plugins::surface` knows about a run of bytes, a name, where it lives and +how it is reached; it does not know what SPIR-V is. `rules-spirv`, `rules-slang` +and `tools-embed` all call it, which is what keeps their generated declarations +from drifting. + +It lives in this package's lib root, so a rule package outside this collection +reaches it by depending on `mcpp:plugins` and activating no feature -- the lib +root alone, which is one small module. That works and is the intended path; +whether the generator should become a package of its own is an open question and +not one this version answers. + + **One copy of the bytes.** A generated data header declares a `static` array, so before this every translation unit that included one carried its own copy. Exactly one translation unit -- the generated implementation -- includes them From 0744232a74f6d7f7c807787e0b4f038b9ec3d3e6 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 00:02:13 +0800 Subject: [PATCH 06/20] mcpp.plugins.island: an island's boundary is generated, not written twice `surface` generates the whole interface for a data payload, because an address and a size are all there is to decide. `island` generates what is mechanical about a code island, and only that: the `extern "C"` header its compiler includes -- guards and `__cplusplus` dance included -- and the module the C++ side imports. THE DECLARATION EXISTS ONCE. Without this a project writes it twice, in a header and again wherever the C++ side reaches it, and C language linkage does not mangle: two copies that disagree are one symbol, the link is clean, and each side reads the arguments by its own ABI with no compile error and no link error. That is the copy this removes, and it is the worst one available at this boundary. `export using ::name;` IS WHY THERE IS NO C PARSER HERE. The module re-exports identifiers rather than restating signatures, so the generator needs only the name before the `(` and the header stays the only place a signature is written. Measured on both implementations this package supports before anything was built on it: a consumer that imports the module and never includes the header calls the entry point and links against an implementation compiled by a different driver, under GCC 16.1 and clang 22.1.8. The C++ interface is still the project's. A seam that turns raw pointers into something worth calling is a design decision no generator makes well; this removes the boilerplate around the boundary, not the boundary's design. `emit_module = false` emits the header alone for a project that keeps a seam that includes rather than imports. `tests/island-interface` uses a `.c` island rather than a `.cu`, because the generator does not know what compiler produced the object and a fixture needing a vendor toolkit could only run where that toolkit is published. It has two entry points, since a generator that re-exported only the first would satisfy a one-function fixture, and the CI step refuses a module that restated a signature even though such a module would compile. The README states why both generators live in this package rather than in mcpp: the engine's module carries the protocol, and a generator is a library on top of it. Code inside the engine changes only with an engine release, which is the coupling this version's `device_extensions` work exists to remove. --- .github/workflows/ci.yml | 39 ++++++ README.md | 63 +++++++++ src/plugins.cppm | 153 +++++++++++++++++++++ tests/island-interface/build.mcpp | 31 +++++ tests/island-interface/mcpp.toml | 38 +++++ tests/island-interface/src/kernels/saxpy.c | 13 ++ tests/island-interface/src/main.cpp | 32 +++++ 7 files changed, 369 insertions(+) create mode 100644 tests/island-interface/build.mcpp create mode 100644 tests/island-interface/mcpp.toml create mode 100644 tests/island-interface/src/kernels/saxpy.c create mode 100644 tests/island-interface/src/main.cpp diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 0d2077e..ed197fa 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -198,6 +198,45 @@ jobs: rm -f dup.log restored.log echo "ok: and it builds again with one shader per stem" + # THE OTHER GENERATED INTERFACE, AND THE ONE THAT IS NOT A PAYLOAD. + # + # `mcpp.plugins.island` generates the boundary a device island is reached + # across: the `extern "C"` header its compiler includes, and the module the + # C++ side imports. The entry points are declared once, so the copy that + # can drift is gone -- and at a C-linkage boundary that copy is the worst + # one available, because nothing mangles and two disagreeing declarations + # link cleanly. + # + # The island here is a `.c` rather than a `.cu`: the generator does not + # know what compiler produced the object, and a fixture needing a vendor + # toolkit could only run where that toolkit is published. + - name: island interfaces are generated, not written + working-directory: tests/island-interface + run: | + "$MCPP" build + "$MCPP" run | tee run.log + grep -q '^out\[0\]=6 out\[1\]=12 out\[2\]=18 out\[3\]=24' run.log + grep -q '^all ok' run.log + # Neither artefact is written by hand. + test ! -f include/island_interface.kernels.h + d=target/.build-mcpp/out/island + grep -q 'extern "C" {' "$d/island_interface.kernels.h" + grep -q '^export using ::saxpy_device;' "$d/island_interface.kernels.cppm" + grep -q '^export using ::scale_device;' "$d/island_interface.kernels.cppm" + # The module re-exports NAMES: a second copy of a signature is the + # thing this generator exists to remove, so its appearance here is a + # regression even though it would compile. + if grep -q 'int saxpy_device(' "$d/island_interface.kernels.cppm"; then + echo "FAIL: the module restated a signature instead of re-exporting a name" + exit 1 + fi + # And the consumer includes nothing. + if grep -rn '#include' src/main.cpp; then + echo "FAIL: the C++ side includes rather than imports" + exit 1 + fi + echo "ok: one declaration, two generated artefacts, no restated signature" + # NO BUILD PROGRAM AT ALL, WHICH IS WHAT THE TWO KEYS BUY. # # `spirv-consumer` and `spirv-module-consumer` both write a `build.mcpp`, diff --git a/README.md b/README.md index 75a024c..bc04797 100644 --- a/README.md +++ b/README.md @@ -261,6 +261,69 @@ before this every translation unit that included one carried its own copy. Exactly one translation unit -- the generated implementation -- includes them now, and every consumer reaches the same array through the accessor. +## The other generated interface: an island's boundary + +`mcpp::plugins::surface` generates the whole interface for a **data** payload, +because an address and a size are all there is to decide. +`mcpp::plugins::island` generates what is mechanical about a **code** island -- +and only that. + +```cpp +// build.mcpp +const std::vector entries{ + "int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n)", +}; +mcpp::plugins::island::options opt; +opt.module_name = "myapp.kernels"; +opt.out_dir = std::string(mcpp::out_dir()) + "/island"; +const auto out = mcpp::plugins::island::emit(entries, opt); +mcpp::include_dir(out->include_dir.c_str()); +mcpp::generated(out->interface_file.c_str()); +``` + +Two files come out of that one declaration: the `extern "C"` header the device +translation unit includes, guards and `__cplusplus` dance included, and the +module the C++ side imports. + +**The declaration still exists once.** Without this, a project writes it twice -- +in a header, and again wherever the C++ side reaches it. C language linkage does +not mangle, so two copies that disagree are one symbol: the link is clean and +each side reads the arguments by its own ABI, with no compile error and no link +error. That is the copy this removes. + +**The module re-exports names, not signatures.** `export using ::saxpy_device;` +needs the identifier and nothing else, so the generator has no C parser in it and +the header stays the only place a signature is written. Measured on both +implementations this package supports: a consumer that imports the module and +never includes the header calls the entry point and links against an +implementation built by a different driver, under GCC 16.1 and clang 22.1.8. + +**The C++ interface is still yours.** A seam that turns raw pointers into +`std::optional>` is a design decision, and no generator makes +it well. This removes the boilerplate around the boundary, not the boundary's +design. `emit_module = false` emits the header alone for a project that keeps a +hand-written seam that includes rather than imports. + +### Why the generators live here and not in mcpp + +The engine's `mcpp` module is compiled into the mcpp binary and carries the +**protocol** -- what a build program can tell mcpp: `action`, `generated`, +`include_dir`, `fact`, `floor`. A generator is a **library on top of** that +protocol: it reads declarations, writes files, and hands them back through +`mcpp::generated`. It extends nothing. + +So the line is: the engine's module carries the protocol, and generators are +libraries. Both of these are opinionated and will change -- what a `payload` +looks like, how a namespace is derived, which storages exist -- and code inside +the engine changes only with an engine release, which is the coupling this +version's `device_extensions` work exists to remove. + +A rule package outside this collection reaches them by depending on +`mcpp:plugins` and activating no feature: the lib root alone, one small module. +That is a package dependency it chooses, not a coupling the engine imposes. If +they stabilise, moving them into the engine later is a low-risk step; the reverse +is not. + ## How the engine sees this package mcpp compiles every module interface unit among a host-module package's diff --git a/src/plugins.cppm b/src/plugins.cppm index 3bc2bf6..139fb37 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -641,3 +641,156 @@ inline std::string module_root_from_package() { } } // namespace mcpp::plugins::surface + +// mcpp.plugins.island -- the boundary a device island is reached across. +// +// WHAT THIS GENERATES AND WHY IT IS NOT THE SAME THING AS `surface`. +// +// `surface` writes the whole interface for a DATA payload, because an address +// and a size are all there is to decide. An island is CODE, and its C++ +// interface -- which functions, which types, what happens on failure -- is a +// design decision no generator makes well. That interface stays hand-written. +// +// What is mechanical is everything AROUND the entry points: an include guard, +// an `extern "C"` block, the `__cplusplus` dance, and a module wrapper whose +// only content is a global module fragment and a re-export. Ten lines of +// boilerplate per one line of content, written the same way in every project +// that has an island. That is what this generates. +// +// THE DECLARATION STILL EXISTS ONCE. It moves from a hand-written header into +// the build program, and both artefacts are produced from it -- the header the +// device compiler includes and the module the C++ side imports. Splitting a +// declaration across those two is the failure this removes, and it is the worst +// one available at this boundary: C language linkage does not mangle, so two +// copies that disagree are one symbol, the link is clean, and each side reads +// the arguments by its own ABI. +// +// `export using ::name;` IS WHAT MAKES IT WORK WITHOUT A C PARSER. +// +// The module re-exports names rather than restating signatures, so the +// generator needs only the identifier before the `(`. Measured with GCC 16.1: +// a consumer that imports the module and never includes the header calls the +// entry point and links against an implementation compiled by a DIFFERENT +// driver, which is the arrangement a device island actually has. +// +// IT IS OPTIONAL, AND A PROJECT THAT WRITES ITS OWN HEADER KEEPS IT. Nothing +// here is required to have an island; it removes boilerplate from projects that +// want it removed. +export namespace mcpp::plugins::island { + +struct options { + // The module the C++ side imports: `myapp.kernels`. The header is named + // after it too, so one name places both files. + std::string module_name; + // Where they are written. The caller puts this on the include path so the + // device translation unit can include the header. + std::string out_dir; + // For the generated files' first line: "mcpp.rules.cuda". + std::string produced_by; + // False emits the header alone, for a project that wants the boilerplate + // removed but keeps a hand-written seam that includes rather than imports. + bool emit_module = true; +}; + +struct emitted { + std::string header_file; // what the device translation unit includes + std::string interface_file; // the `.cppm`; empty when `emit_module` is false + std::string include_dir; // for `mcpp::include_dir` + std::string module_name; // empty when `emit_module` is false +}; + +// The identifier immediately before the first `(`. That is the whole parse this +// needs: the module re-exports the NAME and the header carries the signature +// verbatim, so nothing here has to understand a C declaration. +inline std::string entry_name(std::string_view decl) { + const auto paren = decl.find('('); + if (paren == std::string_view::npos) return {}; + auto end = paren; + while (end > 0 && (decl[end - 1] == ' ' || decl[end - 1] == '\t')) --end; + auto begin = end; + while (begin > 0) { + const char c = decl[begin - 1]; + if (std::isalnum(static_cast(c)) || c == '_') --begin; + else break; + } + return std::string(decl.substr(begin, end - begin)); +} + +inline std::optional emit(std::span entries, + const options& opt) { + if (entries.empty()) return emitted{}; + if (opt.module_name.empty() || opt.out_dir.empty()) { + std::cerr << "mcpp.plugins.island: module_name and out_dir are required\n"; + return std::nullopt; + } + const auto by = opt.produced_by.empty() ? std::string("mcpp.plugins.island") + : opt.produced_by; + const auto dir = std::filesystem::path(opt.out_dir); + + // A name the generator could not find is refused rather than skipped: a + // declaration that produced no re-export would compile, and the consumer's + // failure would be an unresolved name three files away. + std::vector names; + for (auto const& e : entries) { + auto n = entry_name(e); + if (n.empty()) { + std::cerr << std::format( + "mcpp.plugins.island: cannot find an entry point name in `{}`.\n" + " Each entry is a C declaration, e.g.\n" + " \"int saxpy_device(float a, const float* x, unsigned n)\"\n", e); + return std::nullopt; + } + names.push_back(std::move(n)); + } + + emitted out; + const auto guard = [&] { + std::string g = "MCPP_ISLAND_"; + for (char c : opt.module_name) + g += std::isalnum(static_cast(c)) + ? static_cast(std::toupper(static_cast(c))) : '_'; + return g + "_H"; + }(); + + std::string h; + h += std::format("// Generated by mcpp.plugins.island for {0}. Do not edit.\n" + "//\n" + "// The island's boundary. Included by the device translation unit, which\n" + "// is compiled by a compiler mcpp did not resolve -- so the interface is\n" + "// `extern \"C\"`, because the two sides share no C++ ABI.\n" + "#ifndef {1}\n#define {1}\n" + "#ifdef __cplusplus\nextern \"C\" {{\n#endif\n\n", by, guard); + for (auto const& e : entries) h += e + ";\n"; + h += "\n#ifdef __cplusplus\n}\n#endif\n#endif\n"; + + out.header_file = (dir / (opt.module_name + ".h")).string(); + out.include_dir = dir.string(); + if (!surface::write_if_different(out.header_file, h)) { + std::cerr << std::format("mcpp.plugins.island: cannot write {}\n", out.header_file); + return std::nullopt; + } + if (!opt.emit_module) return out; + + std::string m; + m += std::format("// Generated by mcpp.plugins.island for {0}. Do not edit.\n" + "//\n" + "// The C++ side imports this instead of including the header. Names are\n" + "// re-exported rather than restated, so this file carries no second copy\n" + "// of a signature -- which at a C-linkage boundary is the copy that can\n" + "// disagree without anything noticing.\n" + "module;\n#include \"{1}\"\n" + "export module {2};\n\n", by, + std::filesystem::path(out.header_file).filename().string(), + opt.module_name); + for (auto const& n : names) m += std::format("export using ::{};\n", n); + + out.interface_file = (dir / (opt.module_name + ".cppm")).string(); + out.module_name = opt.module_name; + if (!surface::write_if_different(out.interface_file, m)) { + std::cerr << std::format("mcpp.plugins.island: cannot write {}\n", out.interface_file); + return std::nullopt; + } + return out; +} + +} // namespace mcpp::plugins::island diff --git a/tests/island-interface/build.mcpp b/tests/island-interface/build.mcpp new file mode 100644 index 0000000..5b7081b --- /dev/null +++ b/tests/island-interface/build.mcpp @@ -0,0 +1,31 @@ +import std; +import mcpp; +import mcpp.plugins; + +// THE ENTRY POINTS ARE DECLARED ONCE, HERE, AND BOTH ARTEFACTS COME FROM THEM. +// +// Without this the project writes the same declarations twice: in a header for +// the island's compiler, and again in whatever the C++ side reaches them +// through. C language linkage does not mangle, so two copies that disagree are +// one symbol -- the link is clean and each side reads the arguments by its own +// ABI. Generating both from one list removes the copy that can drift. +int main() { + const std::vector entries{ + "int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n)", + "int scale_device(float a, float* out, unsigned n)", + }; + + mcpp::plugins::island::options opt; + opt.module_name = "island_interface.kernels"; + opt.out_dir = std::string(mcpp::out_dir()) + "/island"; + opt.produced_by = "the island-interface fixture"; + + const auto out = mcpp::plugins::island::emit(entries, opt); + if (!out) return 1; + + // The header goes on the include path so the island can include it, and + // both generated files join the build. + mcpp::include_dir(out->include_dir.c_str()); + mcpp::generated(out->interface_file.c_str()); + return 0; +} diff --git a/tests/island-interface/mcpp.toml b/tests/island-interface/mcpp.toml new file mode 100644 index 0000000..c9c8c9b --- /dev/null +++ b/tests/island-interface/mcpp.toml @@ -0,0 +1,38 @@ +# The consumer that tests `mcpp.plugins.island`. +# +# WHY THE ISLAND HERE IS A `.c` AND NOT A `.cu`. +# +# The generator does not know what compiler produced the object. What it knows +# is that the boundary is `extern "C"` and that one side is a module -- which is +# exactly the arrangement a CUDA, HIP, SYCL or Ascend C island has, and is +# reproducible without a vendor toolkit. A fixture that needed one could only +# run where that toolkit is published, and the half of this that could break is +# the generated files rather than the device compiler. +# +# What it asserts: +# +# - the header the island includes is GENERATED, so the guard, the +# `extern "C"` block and the `__cplusplus` dance are not written by hand; +# - the module the C++ side imports is generated too, and re-exports the entry +# points by name rather than restating their signatures; +# - the consumer never includes the header. +[package] +name = "island-interface" +namespace = "example" +version = "0.1.0" +description = "An extern \"C\" island reached through a generated module" + +[language] +standard = "c++23" +modules = true +import_std = true + +[build-dependencies.mcpp] +plugins = { path = "../..", host-module = true } + +[build] +sources = ["src/*.cpp", "src/kernels/*.c"] + +[targets.island-interface] +kind = "bin" +main = "src/main.cpp" diff --git a/tests/island-interface/src/kernels/saxpy.c b/tests/island-interface/src/kernels/saxpy.c new file mode 100644 index 0000000..a6369cf --- /dev/null +++ b/tests/island-interface/src/kernels/saxpy.c @@ -0,0 +1,13 @@ +/* The island. It includes the GENERATED header, which is the only place the + entry points' signatures exist. */ +#include "island_interface.kernels.h" + +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; +} + +int scale_device(float a, float* out, unsigned n) { + for (unsigned i = 0; i < n; ++i) out[i] = a * out[i]; + return 0; +} diff --git a/tests/island-interface/src/main.cpp b/tests/island-interface/src/main.cpp new file mode 100644 index 0000000..7170aea --- /dev/null +++ b/tests/island-interface/src/main.cpp @@ -0,0 +1,32 @@ +// The C++ side. It imports the generated module and includes nothing: the +// header exists for the island's compiler, which does not read modules. +import std; +import island_interface.kernels; + +int main() { + float x[4] = {1, 2, 3, 4}; + float y[4] = {10, 20, 30, 40}; + float out[4] = {}; + + if (saxpy_device(2.0f, x, y, out, 4) != 0) { + std::cout << "BAD: saxpy_device failed\n"; + return 1; + } + if (scale_device(0.5f, out, 4) != 0) { + std::cout << "BAD: scale_device failed\n"; + return 1; + } + + // (2*1+10)/2, (2*2+20)/2, (2*3+30)/2, (2*4+40)/2 + const float want[4] = {6, 12, 18, 24}; + bool ok = true; + for (int i = 0; i < 4; ++i) { + std::cout << std::format("out[{}]={} ", i, out[i]); + if (out[i] != want[i]) ok = false; + } + std::cout << "\n"; + // Two entry points rather than one, because a generator that re-exported + // only the first would still satisfy a single-function fixture. + std::cout << (ok ? "all ok\n" : "FAILED\n"); + return ok ? 0 : 1; +} From 9ef04f6e99b66c3b98d00314faa26abd76628540 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 00:08:39 +0800 Subject: [PATCH 07/20] island: the entry points are marked where they are defined, and the marker names the mechanism `island::scan` reads the marked declarations out of the island's own source, so the signature sits beside the definition and exists once. `island::emit` still takes a list directly, for entry points a scan cannot see, and a project that wants neither writes its own header and module wrapper. Each layer overrides the one above; the default is the one that keeps the signature in one place. THE MARKER SELECTS, AND THAT IS THE POINT. An island has internal functions, and a generator that exported whatever the file contained would make the boundary an accident of the file's contents. The fixture has a static helper and an unmarked external function, and CI refuses either reaching the generated files -- a check that would otherwise pass, because both files would merely be larger. `MCPP_EXPORT_C` names the mechanism rather than the domain: what is marked is exported across a generated boundary with C linkage, and both halves are in the name. `MCPP_ISLAND_EXPORT` was precise inside docs/20's vocabulary and narrow outside it -- the generator works for any `extern "C"` boundary, and the fixture proves it with a plain `.c`. Anything ending `_API` was rejected outright: that suffix conventionally expands to a visibility attribute and this expands to nothing, so borrowing it would promise what it does not do and collide with a project that later wants the real thing. THE SCAN IS NOT A C PARSER. From the marker it copies verbatim to the parenthesis that closes the parameter list, matching nesting. The fixture's second entry point wraps across three lines for that reason, and CI asserts the flattened declaration, because a line-oriented scan would pass every other check here. A GCC 16 defect was found and worked around while writing it: `for (char c : s)` over a `std::string` inside an exported inline function makes the consumer's build program fail with `inlining failed in call to 'always_inline' ... function body not available` in ``, naming neither this file nor the loop. Indexing touches no iterator type. The comment at the site records the message so the next person recognises it. The island still includes the generated header, and the README says why that line is not new: a project writes the same include today and writes the header as well. It also earns its place twice, defining the marker and declaring the entry points, so a definition whose signature drifted fails where it was written rather than at the link. Removing it with `-include` is left as an option, since a file whose marker comes from a flag is no longer self-contained. --- .github/workflows/ci.yml | 20 +++- README.md | 55 ++++++++-- src/plugins.cppm | 120 ++++++++++++++++++++- tests/island-interface/build.mcpp | 36 ++++--- tests/island-interface/src/kernels/saxpy.c | 23 +++- 5 files changed, 230 insertions(+), 24 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index ed197fa..5ccafe6 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -235,7 +235,25 @@ jobs: echo "FAIL: the C++ side includes rather than imports" exit 1 fi - echo "ok: one declaration, two generated artefacts, no restated signature" + # THE MARKER SELECTS, AND THE UNMARKED MUST NOT TRAVEL. The island has + # a static helper and an unmarked external function; a scan that + # exported whatever the file contained would still pass every check + # above, because both generated files would merely be larger. + for name in internal_device unexported_helper; do + if grep -q "$name" "$d/island_interface.kernels.h" \ + || grep -q "$name" "$d/island_interface.kernels.cppm"; then + echo "FAIL: $name is not marked and reached the generated boundary" + exit 1 + fi + done + # The marker is defined by the header, so the island needs no other + # arrangement to compile. + grep -q '^#define MCPP_EXPORT_C$' "$d/island_interface.kernels.h" + # A signature that wrapped across lines survived: the scan matches + # parentheses rather than reading lines. + grep -q 'int scale_device(float a, float\* out, unsigned n);' \ + "$d/island_interface.kernels.h" + echo "ok: one declaration, two generated artefacts, and only what was marked" # NO BUILD PROGRAM AT ALL, WHICH IS WHAT THE TWO KEYS BUY. # diff --git a/README.md b/README.md index bc04797..fcf8fb7 100644 --- a/README.md +++ b/README.md @@ -268,22 +268,63 @@ because an address and a size are all there is to decide. `mcpp::plugins::island` generates what is mechanical about a **code** island -- and only that. +The entry points are marked where they are defined, and the signature exists +once: + +```c +// src/kernels/saxpy.cu +#include "myapp.kernels.h" // generated; defines the marker as nothing + +MCPP_EXPORT_C +int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n) { ... } +``` + ```cpp // build.mcpp -const std::vector entries{ - "int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n)", -}; mcpp::plugins::island::options opt; opt.module_name = "myapp.kernels"; opt.out_dir = std::string(mcpp::out_dir()) + "/island"; -const auto out = mcpp::plugins::island::emit(entries, opt); + +const auto entries = mcpp::plugins::island::scan(islands, opt); +const auto out = mcpp::plugins::island::emit(*entries, opt); mcpp::include_dir(out->include_dir.c_str()); mcpp::generated(out->interface_file.c_str()); ``` -Two files come out of that one declaration: the `extern "C"` header the device -translation unit includes, guards and `__cplusplus` dance included, and the -module the C++ side imports. +Two files come out of that one marked declaration: the `extern "C"` header the +device translation unit includes, guards and `__cplusplus` dance included, and +the module the C++ side imports. + +**Three layers, and each overrides the one above.** `scan` reads the marked +declarations out of the island, which puts the signature beside the definition; +`emit` takes a list directly, for entry points a scan cannot see; and a project +that wants neither writes its own header and its own module wrapper. The default +is the one that keeps the signature in one place. + +**The marker selects.** An island has internal functions, and a generator that +exported whatever the file contained would make the boundary an accident of the +file's contents. `MCPP_EXPORT_C` names the mechanism rather than the domain -- +what is marked is exported across a generated boundary with C linkage -- and +deliberately does not end in `_API`, a suffix that conventionally expands to a +visibility attribute where this expands to nothing. It is configurable through +`options::marker`. + +**The scan is not a C parser.** From the marker it copies verbatim to the +parenthesis closing the parameter list, matching nesting, so a signature that +wraps across lines or carries a macro travels through unexamined. + +**The island includes the generated header, and that line is not new.** Today a +project writes `#include "saxpy/saxpy.h"` *and* writes that header; here the +include stays and the header is what became generated. It earns its place twice +over: it defines the marker, and it declares the entry points, so a definition +whose signature drifted from its declaration fails where it was written rather +than at the link. + +A rule can remove it by passing `-include ` to the device +compiler, which every compiler this package drives accepts. That is left as an +option rather than made the default: a file whose marker comes from a flag is no +longer self-contained, and opening it says nothing about where `MCPP_EXPORT_C` +came from. **The declaration still exists once.** Without this, a project writes it twice -- in a header, and again wherever the C++ side reaches it. C language linkage does diff --git a/src/plugins.cppm b/src/plugins.cppm index 139fb37..2ba09a6 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -690,6 +690,28 @@ struct options { // False emits the header alone, for a project that wants the boilerplate // removed but keeps a hand-written seam that includes rather than imports. bool emit_module = true; + // The marker `scan()` looks for. It is defined as nothing by the generated + // header, so the island includes that header and then writes the marker in + // front of each entry point it exports. + // + // THE NAME SAYS THE MECHANISM, NOT THE DOMAIN. What is marked is exported + // across a generated boundary, with C linkage; both halves are in the name + // and neither narrows it. Two alternatives were considered: + // + // - `MCPP_ISLAND_EXPORT` is precise inside docs/20's vocabulary and + // narrow outside it. The generator does not know what compiler produced + // the object, and works for any `extern "C"` boundary -- a C library + // shim has one and is not an island. + // - anything ending `_API` was rejected outright. That suffix + // conventionally expands to a visibility attribute + // (`__declspec(dllexport)`, `visibility("default")`), and this expands + // to nothing. Borrowing it would promise something it does not do, and + // would collide with a project that later wants the real thing. + // + // `MCPP__` also leaves room: a future marker read by a + // different generator joins the family rather than inventing a second + // shape. + std::string marker = "MCPP_EXPORT_C"; }; struct emitted { @@ -716,6 +738,97 @@ inline std::string entry_name(std::string_view decl) { return std::string(decl.substr(begin, end - begin)); } +// THE ENTRY POINTS, TAKEN FROM WHERE THEY ARE DEFINED. +// +// `emit` takes a list of declarations, which is exact and requires the project +// to write each signature in its build program. This reads them out of the +// island instead, so the signature lives beside the definition and exists once +// -- which is the arrangement a reader expects and the one that cannot drift. +// +// IT IS NOT A C PARSER, AND DOES NOT NEED TO BE. From the marker it copies +// verbatim up to the parenthesis that closes the parameter list, matching +// nesting so a function pointer parameter does not end it early. What it copies +// is what the header will contain, so anything the island's compiler accepts in +// a declaration -- a macro, a qualifier, a multi-line signature -- travels +// through unexamined. +// +// A marker with no `(` after it is refused rather than skipped: a marked entry +// point that produced no declaration would leave the island defining a function +// nothing declares, and the consumer's failure would be an unresolved name in a +// different file. +inline std::optional> +scan(std::span sources, const options& opt) { + std::vector entries; + for (auto const& src : sources) { + std::ifstream in(src); + if (!in) { + std::cerr << std::format("mcpp.plugins.island: cannot read {}\n", src); + return std::nullopt; + } + std::string text((std::istreambuf_iterator(in)), + std::istreambuf_iterator()); + for (std::size_t at = text.find(opt.marker); at != std::string::npos; + at = text.find(opt.marker, at + 1)) { + // The marker's own definition in the generated header is not an + // entry point. Skipped by requiring a `(` before the next `;` or + // `{`, which a `#define` line does not have. + std::size_t i = at + opt.marker.size(); + int depth = 0; + bool sawOpen = false; + std::size_t end = std::string::npos; + for (; i < text.size(); ++i) { + const char c = text[i]; + if (c == '(') { ++depth; sawOpen = true; } + else if (c == ')') { + if (--depth == 0) { end = i; break; } + } else if (!sawOpen && (c == ';' || c == '{' || c == '\n')) { + if (c == '\n') continue; // a signature may wrap + break; // `;` or `{` with no `(` + } + } + if (end == std::string::npos) { + if (!sawOpen) continue; // a `#define` of the marker + std::cerr << std::format( + "mcpp.plugins.island: {} carries `{}` whose parameter list does not " + "close.\n A marked entry point is one declaration, and the generator " + "copies it verbatim.\n", src, opt.marker); + return std::nullopt; + } + auto decl = text.substr(at + opt.marker.size(), + end + 1 - (at + opt.marker.size())); + // Collapse the runs of whitespace a wrapped signature carries, so + // the header reads as one declaration per line. + // INDEXED RATHER THAN A RANGE-FOR, AND THAT IS NOT A STYLE CHOICE. + // + // `for (char c : decl)` over a `std::string` inside an exported + // inline function makes GCC 16 instantiate `std::string::iterator` + // in this BMI, and the consumer's build program then fails to + // compile with + // + // error: inlining failed in call to 'always_inline' + // __normal_iterator>::operator*(): + // function body not available + // + // in ``, naming neither this file nor this + // loop. Indexing touches no iterator type and compiles. + std::string flat; + bool space = false; + for (std::size_t k = 0; k < decl.size(); ++k) { + const char c = decl[k]; + if (c == '\n' || c == '\t' || c == '\r' || c == ' ') { + if (!flat.empty()) space = true; + } else { + if (space) flat += ' '; + space = false; + flat += c; + } + } + if (!flat.empty()) entries.push_back(std::move(flat)); + } + } + return entries; +} + inline std::optional emit(std::span entries, const options& opt) { if (entries.empty()) return emitted{}; @@ -759,7 +872,12 @@ inline std::optional emit(std::span entries, "// is compiled by a compiler mcpp did not resolve -- so the interface is\n" "// `extern \"C\"`, because the two sides share no C++ ABI.\n" "#ifndef {1}\n#define {1}\n" - "#ifdef __cplusplus\nextern \"C\" {{\n#endif\n\n", by, guard); + "// Defined as nothing so the island can mark its entry points and\n" + "// still compile: the marker is for the generator to find, not for\n" + "// the compiler to act on.\n" + "#ifndef {2}\n#define {2}\n#endif\n" + "#ifdef __cplusplus\nextern \"C\" {{\n#endif\n\n", + by, guard, opt.marker); for (auto const& e : entries) h += e + ";\n"; h += "\n#ifdef __cplusplus\n}\n#endif\n#endif\n"; diff --git a/tests/island-interface/build.mcpp b/tests/island-interface/build.mcpp index 5b7081b..142b808 100644 --- a/tests/island-interface/build.mcpp +++ b/tests/island-interface/build.mcpp @@ -2,29 +2,41 @@ import std; import mcpp; import mcpp.plugins; -// THE ENTRY POINTS ARE DECLARED ONCE, HERE, AND BOTH ARTEFACTS COME FROM THEM. +// THE ENTRY POINTS ARE READ FROM WHERE THEY ARE DEFINED. // -// Without this the project writes the same declarations twice: in a header for -// the island's compiler, and again in whatever the C++ side reaches them -// through. C language linkage does not mangle, so two copies that disagree are -// one symbol -- the link is clean and each side reads the arguments by its own -// ABI. Generating both from one list removes the copy that can drift. +// `island::scan` takes them out of the island's own source, so the signature +// sits beside the definition and exists once. The alternative -- and what +// `island::emit` takes directly -- is a list written here, which is exact but +// puts the signatures somewhere other than the code they describe. +// +// Without either, a project writes each declaration twice: in a header for the +// island's compiler and again wherever the C++ side reaches it. C language +// linkage does not mangle, so two copies that disagree are one symbol: the link +// is clean and each side reads the arguments by its own ABI. int main() { - const std::vector entries{ - "int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n)", - "int scale_device(float a, float* out, unsigned n)", + const std::vector islands{ + std::string(mcpp::manifest_dir()) + "/src/kernels/saxpy.c", }; + for (auto const& f : islands) mcpp::rerun_if_changed(f.c_str()); mcpp::plugins::island::options opt; opt.module_name = "island_interface.kernels"; opt.out_dir = std::string(mcpp::out_dir()) + "/island"; opt.produced_by = "the island-interface fixture"; - const auto out = mcpp::plugins::island::emit(entries, opt); + const auto entries = mcpp::plugins::island::scan(islands, opt); + if (!entries) return 1; + if (entries->size() != 2) { + std::cerr << std::format("expected two marked entry points, found {}\n", + entries->size()); + return 1; + } + + const auto out = mcpp::plugins::island::emit(*entries, opt); if (!out) return 1; - // The header goes on the include path so the island can include it, and - // both generated files join the build. + // The header goes on the include path so the island can include it, and the + // module joins the build. mcpp::include_dir(out->include_dir.c_str()); mcpp::generated(out->interface_file.c_str()); return 0; diff --git a/tests/island-interface/src/kernels/saxpy.c b/tests/island-interface/src/kernels/saxpy.c index a6369cf..6f91bee 100644 --- a/tests/island-interface/src/kernels/saxpy.c +++ b/tests/island-interface/src/kernels/saxpy.c @@ -1,13 +1,30 @@ -/* The island. It includes the GENERATED header, which is the only place the - entry points' signatures exist. */ +/* The island. + * + * The signatures live HERE, beside the definitions, and exist once. The marker + * is what the generator finds them by; the generated header defines it as + * nothing, so this file compiles with no other arrangement. + * + * `scale_device` deliberately wraps across lines: a signature that did not fit + * on one is the shape a line-oriented scan gets wrong, and the generator + * matches parentheses rather than reading lines. */ #include "island_interface.kernels.h" +MCPP_EXPORT_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; } -int scale_device(float a, float* out, unsigned n) { +MCPP_EXPORT_C +int scale_device(float a, + float* out, + unsigned n) { for (unsigned i = 0; i < n; ++i) out[i] = a * out[i]; return 0; } + +/* Not marked, so it must not appear in the header or the module: an island has + * internal functions, and exporting them all would make the boundary whatever + * the file happened to contain. */ +static int unexported_helper(int x) { return x + 1; } +int internal_device(int x) { return unexported_helper(x); } From d8ff7d950633f2ad7cf85084c62989bdb06182ef Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 00:14:25 +0800 Subject: [PATCH 08/20] island: the island writes no include either, and the check that line did survives `force_include_flags` returns what makes a compiler read the generated header before the island's first line -- `-include ` for gcc and clang, `/FI` for MSVC. A project using this generator now has no header in its source tree and no line naming one; the fixture's island is a `.c` with zero includes. The line was defended in the previous commit on the grounds that it kept the file self-contained. It did not. The header is generated, so an `#include` of it names a file that does not exist until mcpp writes it, and that translation unit could not be compiled outside mcpp with or without the line. What the line actually bought was a name its author never opens. THE CHECK IT ALSO DID IS STILL THERE, and this was verified rather than assumed. The compiler still sees the declarations, so a definition whose signature drifted fails where it was written: src/kernels/saxpy.c:22:5: error: conflicting types for 'scale_device' That is only reachable on the `emit` path, where a hand-written list and the definition are separate things. Under `scan` the header is generated FROM the definition and the two cannot disagree, so the check is vacuous there -- worth saying, because a criterion that cannot fail is one this project has paid for before. The first attempt at that control reported success for the wrong reason: the fixture was copied to a scratch directory, its `path = "../.."` dependency stopped resolving, and the build failed with `has no mcpp.toml` while the check read "refused". Re-run in place, it fails on the conflicting type. --- .github/workflows/ci.yml | 17 +++++++++-- README.md | 35 ++++++++++++++-------- src/plugins.cppm | 26 +++++++++++++++- tests/island-interface/build.mcpp | 26 ++++++++++------ tests/island-interface/src/kernels/saxpy.c | 11 ++++--- 5 files changed, 86 insertions(+), 29 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 5ccafe6..825bc81 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -230,11 +230,22 @@ jobs: echo "FAIL: the module restated a signature instead of re-exporting a name" exit 1 fi - # And the consumer includes nothing. - if grep -rn '#include' src/main.cpp; then - echo "FAIL: the C++ side includes rather than imports" + # NEITHER SIDE INCLUDES ANYTHING. The C++ side imports, and the + # island reads the generated header through the compiler's + # forced-include flag -- so a project using this generator has no + # header in its source tree and no line naming one. + if grep -rn '#include' src/main.cpp src/kernels/saxpy.c; then + echo "FAIL: a source names an include; the generator exists to remove it" exit 1 fi + # AND THE CHECK THAT INCLUDE USED TO DO IS STILL THERE. The compiler + # sees the declarations, so a definition whose signature drifted fails + # where it was written. Verified by hand against this fixture with a + # hand-written entry list declaring `double n` where the definition + # says `unsigned n`: `saxpy.c:22: error: conflicting types for + # 'scale_device'`, at the definition rather than at the link. + grep -q 'include' build.mcpp \ + || { echo "FAIL: the build program no longer forces the header in"; exit 1; } # THE MARKER SELECTS, AND THE UNMARKED MUST NOT TRAVEL. The island has # a static helper and an unmarked external function; a scan that # exported whatever the file contained would still pass every check diff --git a/README.md b/README.md index fcf8fb7..6db1c7b 100644 --- a/README.md +++ b/README.md @@ -313,18 +313,29 @@ visibility attribute where this expands to nothing. It is configurable through parenthesis closing the parameter list, matching nesting, so a signature that wraps across lines or carries a macro travels through unexamined. -**The island includes the generated header, and that line is not new.** Today a -project writes `#include "saxpy/saxpy.h"` *and* writes that header; here the -include stays and the header is what became generated. It earns its place twice -over: it defines the marker, and it declares the entry points, so a definition -whose signature drifted from its declaration fails where it was written rather -than at the link. - -A rule can remove it by passing `-include ` to the device -compiler, which every compiler this package drives accepts. That is left as an -option rather than made the default: a file whose marker comes from a flag is no -longer self-contained, and opening it says nothing about where `MCPP_EXPORT_C` -came from. +**The island writes no `#include` either.** `force_include_flags` returns the +flags that make the compiler read the generated header before the island's first +line -- `-include ` for gcc and clang, `/FI` for MSVC -- so a project +using this generator has no header in its source tree and no line naming one. + +An `#include` of the generated header would name a file its author never opens, +and it buys no self-containment: that translation unit could not be compiled +outside mcpp with or without the line, because the file it names does not exist +until mcpp writes it. + +**The check that include used to do is still there.** The compiler sees the +declarations, so a definition whose signature drifted from its declaration fails +where it was written rather than at the link. Verified against the fixture with +a hand-written entry list declaring `double n` where the definition says +`unsigned n`: + +``` +src/kernels/saxpy.c:22:5: error: conflicting types for 'scale_device' +``` + +That check is only reachable on the `emit` path, where the list and the +definition are separate things. Under `scan` the header is generated *from* the +definition, so the two cannot disagree at all. **The declaration still exists once.** Without this, a project writes it twice -- in a header, and again wherever the C++ side reaches it. C language linkage does diff --git a/src/plugins.cppm b/src/plugins.cppm index 2ba09a6..690d1db 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -715,12 +715,36 @@ struct options { }; struct emitted { - std::string header_file; // what the device translation unit includes + std::string header_file; // the generated boundary header std::string interface_file; // the `.cppm`; empty when `emit_module` is false std::string include_dir; // for `mcpp::include_dir` std::string module_name; // empty when `emit_module` is false }; +// The flags that make a compiler read the generated header before the island's +// first line, so the island writes neither an include nor anything else. +// +// WHY THIS IS THE DEFAULT RATHER THAN AN INCLUDE LINE. The header is generated: +// it is not in the source tree, and a project using this generator has no +// hand-written header at all. An `#include` of it is therefore a line naming a +// file its author never opens, and it buys no self-containment -- that `.c` +// could not be compiled outside mcpp with or without it, because the file it +// names does not exist until mcpp writes it. +// +// NOTHING IS LOST BY REMOVING IT. The compiler still sees the declarations, so +// a definition whose signature drifted from its declaration still fails where it +// was written rather than at the link, which is the second job the include did. +// A project that prefers the line keeps it: the header carries a guard, so +// including it as well is a no-op. +// +// `/FI` takes the path as one token and `-include` takes it as two, which is why +// this returns a vector rather than a string. +inline std::vector force_include_flags(const std::string& header, + std::string_view compilerId) { + if (compilerId == "msvc") return { "/FI" + header }; + return { "-include", header }; +} + // The identifier immediately before the first `(`. That is the whole parse this // needs: the module re-exports the NAME and the header carries the signature // verbatim, so nothing here has to understand a C declaration. diff --git a/tests/island-interface/build.mcpp b/tests/island-interface/build.mcpp index 142b808..edd6ae3 100644 --- a/tests/island-interface/build.mcpp +++ b/tests/island-interface/build.mcpp @@ -2,15 +2,17 @@ import std; import mcpp; import mcpp.plugins; -// THE ENTRY POINTS ARE READ FROM WHERE THEY ARE DEFINED. +// THE ENTRY POINTS ARE READ FROM WHERE THEY ARE DEFINED, AND THE ISLAND WRITES +// NOTHING ELSE. // -// `island::scan` takes them out of the island's own source, so the signature -// sits beside the definition and exists once. The alternative -- and what -// `island::emit` takes directly -- is a list written here, which is exact but -// puts the signatures somewhere other than the code they describe. +// `island::scan` takes the marked declarations out of the island's own source, +// so the signature sits beside the definition and exists once. `island::emit` +// then writes the boundary header and the module, and the forced-include flags +// below make the island's compiler read that header before the island's first +// line -- so the island has no `#include` either. // -// Without either, a project writes each declaration twice: in a header for the -// island's compiler and again wherever the C++ side reaches it. C language +// Without any of this a project writes each declaration twice: in a header for +// the island's compiler and again wherever the C++ side reaches it. C language // linkage does not mangle, so two copies that disagree are one symbol: the link // is clean and each side reads the arguments by its own ABI. int main() { @@ -35,8 +37,14 @@ int main() { const auto out = mcpp::plugins::island::emit(*entries, opt); if (!out) return 1; - // The header goes on the include path so the island can include it, and the - // module joins the build. + // The C island reads the generated header through `-include`, so the marker + // is defined and the declarations are visible -- which is what keeps a + // signature that drifted from its definition an error where it was written + // rather than at the link. + for (auto const& f : mcpp::plugins::island::force_include_flags( + out->header_file, mcpp::compiler())) + mcpp::cflag(f.c_str()); + mcpp::include_dir(out->include_dir.c_str()); mcpp::generated(out->interface_file.c_str()); return 0; diff --git a/tests/island-interface/src/kernels/saxpy.c b/tests/island-interface/src/kernels/saxpy.c index 6f91bee..fe55b0a 100644 --- a/tests/island-interface/src/kernels/saxpy.c +++ b/tests/island-interface/src/kernels/saxpy.c @@ -1,13 +1,16 @@ -/* The island. +/* The island. No include, and no header in this project at all. * * The signatures live HERE, beside the definitions, and exist once. The marker - * is what the generator finds them by; the generated header defines it as - * nothing, so this file compiles with no other arrangement. + * is what the generator finds them by, and the generated boundary header + * reaches this file through the compiler's forced-include flag rather than + * through a line naming a file that is not in the source tree. + * + * That header still declares these functions, so a definition whose signature + * drifted from its declaration fails here rather than at the link. * * `scale_device` deliberately wraps across lines: a signature that did not fit * on one is the shape a line-oriented scan gets wrong, and the generator * matches parentheses rather than reading lines. */ -#include "island_interface.kernels.h" MCPP_EXPORT_C int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n) { From 7b4cfbc3e218f6473d2f47903a44d0dbe183a023 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 00:21:05 +0800 Subject: [PATCH 09/20] tools: the island generator is a member, not part of the lib root The lib root carries what members share. `mcpp::plugins::surface` is there because three rules produce the same declarations for a consumer and a fourth copy would drift. Nothing in this collection uses the island generator: a project calls it from its own `build.mcpp`, exactly as it calls `mcpp.tools.embed`. Leaving it in the lib root handed it to every consumer whether or not they asked for it, while `tools-embed` beside it required a feature. It moves to `tools/island.cppm` behind `tools-island`, exporting `mcpp.tools.island`. `all-rules-compile` activates the feature and calls `entry_name`, so the parser is compiled and exercised on every host. --- .github/workflows/ci.yml | 2 +- README.md | 14 +- mcpp.toml | 7 +- src/plugins.cppm | 295 --------------------------- tests/all-rules-compile/build.mcpp | 4 + tests/all-rules-compile/mcpp.toml | 2 +- tests/island-interface/build.mcpp | 10 +- tests/island-interface/mcpp.toml | 4 +- tools/island.cppm | 317 +++++++++++++++++++++++++++++ 9 files changed, 343 insertions(+), 312 deletions(-) create mode 100644 tools/island.cppm diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 825bc81..07ff331 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -200,7 +200,7 @@ jobs: # THE OTHER GENERATED INTERFACE, AND THE ONE THAT IS NOT A PAYLOAD. # - # `mcpp.plugins.island` generates the boundary a device island is reached + # `mcpp.tools.island` generates the boundary a device island is reached # across: the `extern "C"` header its compiler includes, and the module the # C++ side imports. The entry points are declared once, so the copy that # can drift is gone -- and at a C-linkage boundary that copy is the worst diff --git a/README.md b/README.md index 6db1c7b..09f92e5 100644 --- a/README.md +++ b/README.md @@ -53,6 +53,7 @@ engine's own module family and is not used here. | `rules-spirv` | `mcpp.rules.spirv` | 2026.9.6.6 | `[build] accel = "vulkan1.2"`, a constrained glob for the shader stages; compiles each shader through a `role = "source"` action and states which of the two compilers produced it | | `rules-sycl` | `mcpp.rules.sycl` | 2026.9.6.6 | `[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. Its own engine need is `.sycl` in the device-source table, 2026.9.6.1 | | `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 | +| `tools-island` | `mcpp.tools.island` | 2026.9.7.1 | nothing beyond mcpp: it reads marked entry points out of an island's own source and writes the `extern "C"` boundary header its compiler reads and the module the C++ side imports. Not a device rule -- it claims no extension, and a project calls it from its own `build.mcpp` | ### Each rule brings its own environment @@ -261,12 +262,13 @@ before this every translation unit that included one carried its own copy. Exactly one translation unit -- the generated implementation -- includes them now, and every consumer reaches the same array through the accessor. -## The other generated interface: an island's boundary +## `tools-island`: an island's boundary `mcpp::plugins::surface` generates the whole interface for a **data** payload, because an address and a size are all there is to decide. -`mcpp::plugins::island` generates what is mechanical about a **code** island -- -and only that. +`mcpp.tools.island` generates what is mechanical about a **code** island -- and +only that. It is a member like `tools-embed` rather than part of the lib root: +nothing in this collection uses it, a project does. The entry points are marked where they are defined, and the signature exists once: @@ -281,12 +283,12 @@ int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n ```cpp // build.mcpp -mcpp::plugins::island::options opt; +mcpp::tools::island::options opt; opt.module_name = "myapp.kernels"; opt.out_dir = std::string(mcpp::out_dir()) + "/island"; -const auto entries = mcpp::plugins::island::scan(islands, opt); -const auto out = mcpp::plugins::island::emit(*entries, opt); +const auto entries = mcpp::tools::island::scan(islands, opt); +const auto out = mcpp::tools::island::emit(*entries, opt); mcpp::include_dir(out->include_dir.c_str()); mcpp::generated(out->interface_file.c_str()); ``` diff --git a/mcpp.toml b/mcpp.toml index ec50185..85c431e 100644 --- a/mcpp.toml +++ b/mcpp.toml @@ -81,12 +81,15 @@ sources = ["rules/sycl.cppm"] rule_module = "mcpp.rules.sycl" device_extensions = [".sycl"] -# NOT a device rule: it embeds files that are already on disk, so it claims no -# extension and names no rule module. A consumer calls it from its own +# NOT device rules: these embed or declare things a project already has, so they +# claim no extension and name no rule module. A consumer calls them from its own # `build.mcpp`. [features.tools-embed] sources = ["tools/embed.cppm"] +[features.tools-island] +sources = ["tools/island.cppm"] + # ── The environment each rule needs (mcpp 2026.9.6.6+) ────────────────────── # # A RULE OWNS THE LIST OF PACKAGES IT DRIVES. It is the code that runs the diff --git a/src/plugins.cppm b/src/plugins.cppm index 690d1db..3bc2bf6 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -641,298 +641,3 @@ inline std::string module_root_from_package() { } } // namespace mcpp::plugins::surface - -// mcpp.plugins.island -- the boundary a device island is reached across. -// -// WHAT THIS GENERATES AND WHY IT IS NOT THE SAME THING AS `surface`. -// -// `surface` writes the whole interface for a DATA payload, because an address -// and a size are all there is to decide. An island is CODE, and its C++ -// interface -- which functions, which types, what happens on failure -- is a -// design decision no generator makes well. That interface stays hand-written. -// -// What is mechanical is everything AROUND the entry points: an include guard, -// an `extern "C"` block, the `__cplusplus` dance, and a module wrapper whose -// only content is a global module fragment and a re-export. Ten lines of -// boilerplate per one line of content, written the same way in every project -// that has an island. That is what this generates. -// -// THE DECLARATION STILL EXISTS ONCE. It moves from a hand-written header into -// the build program, and both artefacts are produced from it -- the header the -// device compiler includes and the module the C++ side imports. Splitting a -// declaration across those two is the failure this removes, and it is the worst -// one available at this boundary: C language linkage does not mangle, so two -// copies that disagree are one symbol, the link is clean, and each side reads -// the arguments by its own ABI. -// -// `export using ::name;` IS WHAT MAKES IT WORK WITHOUT A C PARSER. -// -// The module re-exports names rather than restating signatures, so the -// generator needs only the identifier before the `(`. Measured with GCC 16.1: -// a consumer that imports the module and never includes the header calls the -// entry point and links against an implementation compiled by a DIFFERENT -// driver, which is the arrangement a device island actually has. -// -// IT IS OPTIONAL, AND A PROJECT THAT WRITES ITS OWN HEADER KEEPS IT. Nothing -// here is required to have an island; it removes boilerplate from projects that -// want it removed. -export namespace mcpp::plugins::island { - -struct options { - // The module the C++ side imports: `myapp.kernels`. The header is named - // after it too, so one name places both files. - std::string module_name; - // Where they are written. The caller puts this on the include path so the - // device translation unit can include the header. - std::string out_dir; - // For the generated files' first line: "mcpp.rules.cuda". - std::string produced_by; - // False emits the header alone, for a project that wants the boilerplate - // removed but keeps a hand-written seam that includes rather than imports. - bool emit_module = true; - // The marker `scan()` looks for. It is defined as nothing by the generated - // header, so the island includes that header and then writes the marker in - // front of each entry point it exports. - // - // THE NAME SAYS THE MECHANISM, NOT THE DOMAIN. What is marked is exported - // across a generated boundary, with C linkage; both halves are in the name - // and neither narrows it. Two alternatives were considered: - // - // - `MCPP_ISLAND_EXPORT` is precise inside docs/20's vocabulary and - // narrow outside it. The generator does not know what compiler produced - // the object, and works for any `extern "C"` boundary -- a C library - // shim has one and is not an island. - // - anything ending `_API` was rejected outright. That suffix - // conventionally expands to a visibility attribute - // (`__declspec(dllexport)`, `visibility("default")`), and this expands - // to nothing. Borrowing it would promise something it does not do, and - // would collide with a project that later wants the real thing. - // - // `MCPP__` also leaves room: a future marker read by a - // different generator joins the family rather than inventing a second - // shape. - std::string marker = "MCPP_EXPORT_C"; -}; - -struct emitted { - std::string header_file; // the generated boundary header - std::string interface_file; // the `.cppm`; empty when `emit_module` is false - std::string include_dir; // for `mcpp::include_dir` - std::string module_name; // empty when `emit_module` is false -}; - -// The flags that make a compiler read the generated header before the island's -// first line, so the island writes neither an include nor anything else. -// -// WHY THIS IS THE DEFAULT RATHER THAN AN INCLUDE LINE. The header is generated: -// it is not in the source tree, and a project using this generator has no -// hand-written header at all. An `#include` of it is therefore a line naming a -// file its author never opens, and it buys no self-containment -- that `.c` -// could not be compiled outside mcpp with or without it, because the file it -// names does not exist until mcpp writes it. -// -// NOTHING IS LOST BY REMOVING IT. The compiler still sees the declarations, so -// a definition whose signature drifted from its declaration still fails where it -// was written rather than at the link, which is the second job the include did. -// A project that prefers the line keeps it: the header carries a guard, so -// including it as well is a no-op. -// -// `/FI` takes the path as one token and `-include` takes it as two, which is why -// this returns a vector rather than a string. -inline std::vector force_include_flags(const std::string& header, - std::string_view compilerId) { - if (compilerId == "msvc") return { "/FI" + header }; - return { "-include", header }; -} - -// The identifier immediately before the first `(`. That is the whole parse this -// needs: the module re-exports the NAME and the header carries the signature -// verbatim, so nothing here has to understand a C declaration. -inline std::string entry_name(std::string_view decl) { - const auto paren = decl.find('('); - if (paren == std::string_view::npos) return {}; - auto end = paren; - while (end > 0 && (decl[end - 1] == ' ' || decl[end - 1] == '\t')) --end; - auto begin = end; - while (begin > 0) { - const char c = decl[begin - 1]; - if (std::isalnum(static_cast(c)) || c == '_') --begin; - else break; - } - return std::string(decl.substr(begin, end - begin)); -} - -// THE ENTRY POINTS, TAKEN FROM WHERE THEY ARE DEFINED. -// -// `emit` takes a list of declarations, which is exact and requires the project -// to write each signature in its build program. This reads them out of the -// island instead, so the signature lives beside the definition and exists once -// -- which is the arrangement a reader expects and the one that cannot drift. -// -// IT IS NOT A C PARSER, AND DOES NOT NEED TO BE. From the marker it copies -// verbatim up to the parenthesis that closes the parameter list, matching -// nesting so a function pointer parameter does not end it early. What it copies -// is what the header will contain, so anything the island's compiler accepts in -// a declaration -- a macro, a qualifier, a multi-line signature -- travels -// through unexamined. -// -// A marker with no `(` after it is refused rather than skipped: a marked entry -// point that produced no declaration would leave the island defining a function -// nothing declares, and the consumer's failure would be an unresolved name in a -// different file. -inline std::optional> -scan(std::span sources, const options& opt) { - std::vector entries; - for (auto const& src : sources) { - std::ifstream in(src); - if (!in) { - std::cerr << std::format("mcpp.plugins.island: cannot read {}\n", src); - return std::nullopt; - } - std::string text((std::istreambuf_iterator(in)), - std::istreambuf_iterator()); - for (std::size_t at = text.find(opt.marker); at != std::string::npos; - at = text.find(opt.marker, at + 1)) { - // The marker's own definition in the generated header is not an - // entry point. Skipped by requiring a `(` before the next `;` or - // `{`, which a `#define` line does not have. - std::size_t i = at + opt.marker.size(); - int depth = 0; - bool sawOpen = false; - std::size_t end = std::string::npos; - for (; i < text.size(); ++i) { - const char c = text[i]; - if (c == '(') { ++depth; sawOpen = true; } - else if (c == ')') { - if (--depth == 0) { end = i; break; } - } else if (!sawOpen && (c == ';' || c == '{' || c == '\n')) { - if (c == '\n') continue; // a signature may wrap - break; // `;` or `{` with no `(` - } - } - if (end == std::string::npos) { - if (!sawOpen) continue; // a `#define` of the marker - std::cerr << std::format( - "mcpp.plugins.island: {} carries `{}` whose parameter list does not " - "close.\n A marked entry point is one declaration, and the generator " - "copies it verbatim.\n", src, opt.marker); - return std::nullopt; - } - auto decl = text.substr(at + opt.marker.size(), - end + 1 - (at + opt.marker.size())); - // Collapse the runs of whitespace a wrapped signature carries, so - // the header reads as one declaration per line. - // INDEXED RATHER THAN A RANGE-FOR, AND THAT IS NOT A STYLE CHOICE. - // - // `for (char c : decl)` over a `std::string` inside an exported - // inline function makes GCC 16 instantiate `std::string::iterator` - // in this BMI, and the consumer's build program then fails to - // compile with - // - // error: inlining failed in call to 'always_inline' - // __normal_iterator>::operator*(): - // function body not available - // - // in ``, naming neither this file nor this - // loop. Indexing touches no iterator type and compiles. - std::string flat; - bool space = false; - for (std::size_t k = 0; k < decl.size(); ++k) { - const char c = decl[k]; - if (c == '\n' || c == '\t' || c == '\r' || c == ' ') { - if (!flat.empty()) space = true; - } else { - if (space) flat += ' '; - space = false; - flat += c; - } - } - if (!flat.empty()) entries.push_back(std::move(flat)); - } - } - return entries; -} - -inline std::optional emit(std::span entries, - const options& opt) { - if (entries.empty()) return emitted{}; - if (opt.module_name.empty() || opt.out_dir.empty()) { - std::cerr << "mcpp.plugins.island: module_name and out_dir are required\n"; - return std::nullopt; - } - const auto by = opt.produced_by.empty() ? std::string("mcpp.plugins.island") - : opt.produced_by; - const auto dir = std::filesystem::path(opt.out_dir); - - // A name the generator could not find is refused rather than skipped: a - // declaration that produced no re-export would compile, and the consumer's - // failure would be an unresolved name three files away. - std::vector names; - for (auto const& e : entries) { - auto n = entry_name(e); - if (n.empty()) { - std::cerr << std::format( - "mcpp.plugins.island: cannot find an entry point name in `{}`.\n" - " Each entry is a C declaration, e.g.\n" - " \"int saxpy_device(float a, const float* x, unsigned n)\"\n", e); - return std::nullopt; - } - names.push_back(std::move(n)); - } - - emitted out; - const auto guard = [&] { - std::string g = "MCPP_ISLAND_"; - for (char c : opt.module_name) - g += std::isalnum(static_cast(c)) - ? static_cast(std::toupper(static_cast(c))) : '_'; - return g + "_H"; - }(); - - std::string h; - h += std::format("// Generated by mcpp.plugins.island for {0}. Do not edit.\n" - "//\n" - "// The island's boundary. Included by the device translation unit, which\n" - "// is compiled by a compiler mcpp did not resolve -- so the interface is\n" - "// `extern \"C\"`, because the two sides share no C++ ABI.\n" - "#ifndef {1}\n#define {1}\n" - "// Defined as nothing so the island can mark its entry points and\n" - "// still compile: the marker is for the generator to find, not for\n" - "// the compiler to act on.\n" - "#ifndef {2}\n#define {2}\n#endif\n" - "#ifdef __cplusplus\nextern \"C\" {{\n#endif\n\n", - by, guard, opt.marker); - for (auto const& e : entries) h += e + ";\n"; - h += "\n#ifdef __cplusplus\n}\n#endif\n#endif\n"; - - out.header_file = (dir / (opt.module_name + ".h")).string(); - out.include_dir = dir.string(); - if (!surface::write_if_different(out.header_file, h)) { - std::cerr << std::format("mcpp.plugins.island: cannot write {}\n", out.header_file); - return std::nullopt; - } - if (!opt.emit_module) return out; - - std::string m; - m += std::format("// Generated by mcpp.plugins.island for {0}. Do not edit.\n" - "//\n" - "// The C++ side imports this instead of including the header. Names are\n" - "// re-exported rather than restated, so this file carries no second copy\n" - "// of a signature -- which at a C-linkage boundary is the copy that can\n" - "// disagree without anything noticing.\n" - "module;\n#include \"{1}\"\n" - "export module {2};\n\n", by, - std::filesystem::path(out.header_file).filename().string(), - opt.module_name); - for (auto const& n : names) m += std::format("export using ::{};\n", n); - - out.interface_file = (dir / (opt.module_name + ".cppm")).string(); - out.module_name = opt.module_name; - if (!surface::write_if_different(out.interface_file, m)) { - std::cerr << std::format("mcpp.plugins.island: cannot write {}\n", out.interface_file); - return std::nullopt; - } - return out; -} - -} // namespace mcpp::plugins::island diff --git a/tests/all-rules-compile/build.mcpp b/tests/all-rules-compile/build.mcpp index b838d9e..120dbd1 100644 --- a/tests/all-rules-compile/build.mcpp +++ b/tests/all-rules-compile/build.mcpp @@ -12,6 +12,7 @@ import mcpp.rules.slang; import mcpp.rules.spirv; import mcpp.rules.sycl; import mcpp.tools.embed; +import mcpp.tools.island; int main() { // No accelerator is named, so each of these returns true without looking @@ -28,6 +29,9 @@ int main() { // called, because a fixture whose point is a compilation should not also // produce output. ok = ok && !mcpp::tools::embed::header_path("probe.bin").empty(); + // Same question of the island generator: the one call that writes + // nothing. `entry_name` is where its only parsing lives. + ok = ok && mcpp::tools::island::entry_name("int f(int a)") == "f"; std::println("all-rules-compile: every rule module compiled for this host"); return ok ? 0 : 1; } diff --git a/tests/all-rules-compile/mcpp.toml b/tests/all-rules-compile/mcpp.toml index 2979774..78e8c6b 100644 --- a/tests/all-rules-compile/mcpp.toml +++ b/tests/all-rules-compile/mcpp.toml @@ -38,7 +38,7 @@ import_std = true plugins = { path = "../..", features = [ "rules-ascendc", "rules-cuda", "rules-hip", "rules-slang", "rules-spirv", "rules-sycl", - "tools-embed", + "tools-embed", "tools-island", ], host-module = true } # NO `accel`, and that is the whole design: with none, every rule returns diff --git a/tests/island-interface/build.mcpp b/tests/island-interface/build.mcpp index edd6ae3..b80ac47 100644 --- a/tests/island-interface/build.mcpp +++ b/tests/island-interface/build.mcpp @@ -1,6 +1,6 @@ import std; import mcpp; -import mcpp.plugins; +import mcpp.tools.island; // THE ENTRY POINTS ARE READ FROM WHERE THEY ARE DEFINED, AND THE ISLAND WRITES // NOTHING ELSE. @@ -21,12 +21,12 @@ int main() { }; for (auto const& f : islands) mcpp::rerun_if_changed(f.c_str()); - mcpp::plugins::island::options opt; + mcpp::tools::island::options opt; opt.module_name = "island_interface.kernels"; opt.out_dir = std::string(mcpp::out_dir()) + "/island"; opt.produced_by = "the island-interface fixture"; - const auto entries = mcpp::plugins::island::scan(islands, opt); + const auto entries = mcpp::tools::island::scan(islands, opt); if (!entries) return 1; if (entries->size() != 2) { std::cerr << std::format("expected two marked entry points, found {}\n", @@ -34,14 +34,14 @@ int main() { return 1; } - const auto out = mcpp::plugins::island::emit(*entries, opt); + const auto out = mcpp::tools::island::emit(*entries, opt); if (!out) return 1; // The C island reads the generated header through `-include`, so the marker // is defined and the declarations are visible -- which is what keeps a // signature that drifted from its definition an error where it was written // rather than at the link. - for (auto const& f : mcpp::plugins::island::force_include_flags( + for (auto const& f : mcpp::tools::island::force_include_flags( out->header_file, mcpp::compiler())) mcpp::cflag(f.c_str()); diff --git a/tests/island-interface/mcpp.toml b/tests/island-interface/mcpp.toml index c9c8c9b..0f261c6 100644 --- a/tests/island-interface/mcpp.toml +++ b/tests/island-interface/mcpp.toml @@ -1,4 +1,4 @@ -# The consumer that tests `mcpp.plugins.island`. +# The consumer that tests `mcpp.tools.island`. # # WHY THE ISLAND HERE IS A `.c` AND NOT A `.cu`. # @@ -28,7 +28,7 @@ modules = true import_std = true [build-dependencies.mcpp] -plugins = { path = "../..", host-module = true } +plugins = { path = "../..", features = ["tools-island"], host-module = true } [build] sources = ["src/*.cpp", "src/kernels/*.c"] diff --git a/tools/island.cppm b/tools/island.cppm new file mode 100644 index 0000000..95e9ba1 --- /dev/null +++ b/tools/island.cppm @@ -0,0 +1,317 @@ +// mcpp.tools.island -- see the block comment below. +// +// A MEMBER RATHER THAN PART OF THE LIB ROOT, AND THAT IS A CONSISTENCY FIX. +// +// The lib root carries what MEMBERS share: `mcpp::plugins::surface` is there +// because three rules produce the same declarations for a consumer and a fourth +// copy would drift. Nothing in this collection uses the island generator -- a +// project does, directly from its own `build.mcpp`, exactly as it uses +// `mcpp.tools.embed`. Leaving it in the lib root gave it to every consumer +// whether or not they asked, while `tools-embed` next to it required a feature. +module; +#include + +export module mcpp.tools.island; + +import std; +import mcpp; +// For `write_if_different`, which the lib root owns because every generator in +// this package needs the same "do not touch a file whose content is unchanged" +// rule. +import mcpp.plugins; + + +// mcpp.tools.island -- the boundary a device island is reached across. +// +// WHAT THIS GENERATES AND WHY IT IS NOT THE SAME THING AS `surface`. +// +// `surface` writes the whole interface for a DATA payload, because an address +// and a size are all there is to decide. An island is CODE, and its C++ +// interface -- which functions, which types, what happens on failure -- is a +// design decision no generator makes well. That interface stays hand-written. +// +// What is mechanical is everything AROUND the entry points: an include guard, +// an `extern "C"` block, the `__cplusplus` dance, and a module wrapper whose +// only content is a global module fragment and a re-export. Ten lines of +// boilerplate per one line of content, written the same way in every project +// that has an island. That is what this generates. +// +// THE DECLARATION STILL EXISTS ONCE. It moves from a hand-written header into +// the build program, and both artefacts are produced from it -- the header the +// device compiler includes and the module the C++ side imports. Splitting a +// declaration across those two is the failure this removes, and it is the worst +// one available at this boundary: C language linkage does not mangle, so two +// copies that disagree are one symbol, the link is clean, and each side reads +// the arguments by its own ABI. +// +// `export using ::name;` IS WHAT MAKES IT WORK WITHOUT A C PARSER. +// +// The module re-exports names rather than restating signatures, so the +// generator needs only the identifier before the `(`. Measured with GCC 16.1: +// a consumer that imports the module and never includes the header calls the +// entry point and links against an implementation compiled by a DIFFERENT +// driver, which is the arrangement a device island actually has. +// +// IT IS OPTIONAL, AND A PROJECT THAT WRITES ITS OWN HEADER KEEPS IT. Nothing +// here is required to have an island; it removes boilerplate from projects that +// want it removed. +export namespace mcpp::tools::island { + +struct options { + // The module the C++ side imports: `myapp.kernels`. The header is named + // after it too, so one name places both files. + std::string module_name; + // Where they are written. The caller puts this on the include path so the + // device translation unit can include the header. + std::string out_dir; + // For the generated files' first line: "mcpp.rules.cuda". + std::string produced_by; + // False emits the header alone, for a project that wants the boilerplate + // removed but keeps a hand-written seam that includes rather than imports. + bool emit_module = true; + // The marker `scan()` looks for. It is defined as nothing by the generated + // header, so the island includes that header and then writes the marker in + // front of each entry point it exports. + // + // THE NAME SAYS THE MECHANISM, NOT THE DOMAIN. What is marked is exported + // across a generated boundary, with C linkage; both halves are in the name + // and neither narrows it. Two alternatives were considered: + // + // - `MCPP_ISLAND_EXPORT` is precise inside docs/20's vocabulary and + // narrow outside it. The generator does not know what compiler produced + // the object, and works for any `extern "C"` boundary -- a C library + // shim has one and is not an island. + // - anything ending `_API` was rejected outright. That suffix + // conventionally expands to a visibility attribute + // (`__declspec(dllexport)`, `visibility("default")`), and this expands + // to nothing. Borrowing it would promise something it does not do, and + // would collide with a project that later wants the real thing. + // + // `MCPP__` also leaves room: a future marker read by a + // different generator joins the family rather than inventing a second + // shape. + std::string marker = "MCPP_EXPORT_C"; +}; + +struct emitted { + std::string header_file; // the generated boundary header + std::string interface_file; // the `.cppm`; empty when `emit_module` is false + std::string include_dir; // for `mcpp::include_dir` + std::string module_name; // empty when `emit_module` is false +}; + +// The flags that make a compiler read the generated header before the island's +// first line, so the island writes neither an include nor anything else. +// +// WHY THIS IS THE DEFAULT RATHER THAN AN INCLUDE LINE. The header is generated: +// it is not in the source tree, and a project using this generator has no +// hand-written header at all. An `#include` of it is therefore a line naming a +// file its author never opens, and it buys no self-containment -- that `.c` +// could not be compiled outside mcpp with or without it, because the file it +// names does not exist until mcpp writes it. +// +// NOTHING IS LOST BY REMOVING IT. The compiler still sees the declarations, so +// a definition whose signature drifted from its declaration still fails where it +// was written rather than at the link, which is the second job the include did. +// A project that prefers the line keeps it: the header carries a guard, so +// including it as well is a no-op. +// +// `/FI` takes the path as one token and `-include` takes it as two, which is why +// this returns a vector rather than a string. +inline std::vector force_include_flags(const std::string& header, + std::string_view compilerId) { + if (compilerId == "msvc") return { "/FI" + header }; + return { "-include", header }; +} + +// The identifier immediately before the first `(`. That is the whole parse this +// needs: the module re-exports the NAME and the header carries the signature +// verbatim, so nothing here has to understand a C declaration. +inline std::string entry_name(std::string_view decl) { + const auto paren = decl.find('('); + if (paren == std::string_view::npos) return {}; + auto end = paren; + while (end > 0 && (decl[end - 1] == ' ' || decl[end - 1] == '\t')) --end; + auto begin = end; + while (begin > 0) { + const char c = decl[begin - 1]; + if (std::isalnum(static_cast(c)) || c == '_') --begin; + else break; + } + return std::string(decl.substr(begin, end - begin)); +} + +// THE ENTRY POINTS, TAKEN FROM WHERE THEY ARE DEFINED. +// +// `emit` takes a list of declarations, which is exact and requires the project +// to write each signature in its build program. This reads them out of the +// island instead, so the signature lives beside the definition and exists once +// -- which is the arrangement a reader expects and the one that cannot drift. +// +// IT IS NOT A C PARSER, AND DOES NOT NEED TO BE. From the marker it copies +// verbatim up to the parenthesis that closes the parameter list, matching +// nesting so a function pointer parameter does not end it early. What it copies +// is what the header will contain, so anything the island's compiler accepts in +// a declaration -- a macro, a qualifier, a multi-line signature -- travels +// through unexamined. +// +// A marker with no `(` after it is refused rather than skipped: a marked entry +// point that produced no declaration would leave the island defining a function +// nothing declares, and the consumer's failure would be an unresolved name in a +// different file. +inline std::optional> +scan(std::span sources, const options& opt) { + std::vector entries; + for (auto const& src : sources) { + std::ifstream in(src); + if (!in) { + std::cerr << std::format("mcpp.tools.island: cannot read {}\n", src); + return std::nullopt; + } + std::string text((std::istreambuf_iterator(in)), + std::istreambuf_iterator()); + for (std::size_t at = text.find(opt.marker); at != std::string::npos; + at = text.find(opt.marker, at + 1)) { + // The marker's own definition in the generated header is not an + // entry point. Skipped by requiring a `(` before the next `;` or + // `{`, which a `#define` line does not have. + std::size_t i = at + opt.marker.size(); + int depth = 0; + bool sawOpen = false; + std::size_t end = std::string::npos; + for (; i < text.size(); ++i) { + const char c = text[i]; + if (c == '(') { ++depth; sawOpen = true; } + else if (c == ')') { + if (--depth == 0) { end = i; break; } + } else if (!sawOpen && (c == ';' || c == '{' || c == '\n')) { + if (c == '\n') continue; // a signature may wrap + break; // `;` or `{` with no `(` + } + } + if (end == std::string::npos) { + if (!sawOpen) continue; // a `#define` of the marker + std::cerr << std::format( + "mcpp.tools.island: {} carries `{}` whose parameter list does not " + "close.\n A marked entry point is one declaration, and the generator " + "copies it verbatim.\n", src, opt.marker); + return std::nullopt; + } + auto decl = text.substr(at + opt.marker.size(), + end + 1 - (at + opt.marker.size())); + // Collapse the runs of whitespace a wrapped signature carries, so + // the header reads as one declaration per line. + // INDEXED RATHER THAN A RANGE-FOR, AND THAT IS NOT A STYLE CHOICE. + // + // `for (char c : decl)` over a `std::string` inside an exported + // inline function makes GCC 16 instantiate `std::string::iterator` + // in this BMI, and the consumer's build program then fails to + // compile with + // + // error: inlining failed in call to 'always_inline' + // __normal_iterator>::operator*(): + // function body not available + // + // in ``, naming neither this file nor this + // loop. Indexing touches no iterator type and compiles. + std::string flat; + bool space = false; + for (std::size_t k = 0; k < decl.size(); ++k) { + const char c = decl[k]; + if (c == '\n' || c == '\t' || c == '\r' || c == ' ') { + if (!flat.empty()) space = true; + } else { + if (space) flat += ' '; + space = false; + flat += c; + } + } + if (!flat.empty()) entries.push_back(std::move(flat)); + } + } + return entries; +} + +inline std::optional emit(std::span entries, + const options& opt) { + if (entries.empty()) return emitted{}; + if (opt.module_name.empty() || opt.out_dir.empty()) { + std::cerr << "mcpp.tools.island: module_name and out_dir are required\n"; + return std::nullopt; + } + const auto by = opt.produced_by.empty() ? std::string("mcpp.tools.island") + : opt.produced_by; + const auto dir = std::filesystem::path(opt.out_dir); + + // A name the generator could not find is refused rather than skipped: a + // declaration that produced no re-export would compile, and the consumer's + // failure would be an unresolved name three files away. + std::vector names; + for (auto const& e : entries) { + auto n = entry_name(e); + if (n.empty()) { + std::cerr << std::format( + "mcpp.tools.island: cannot find an entry point name in `{}`.\n" + " Each entry is a C declaration, e.g.\n" + " \"int saxpy_device(float a, const float* x, unsigned n)\"\n", e); + return std::nullopt; + } + names.push_back(std::move(n)); + } + + emitted out; + const auto guard = [&] { + std::string g = "MCPP_ISLAND_"; + for (char c : opt.module_name) + g += std::isalnum(static_cast(c)) + ? static_cast(std::toupper(static_cast(c))) : '_'; + return g + "_H"; + }(); + + std::string h; + h += std::format("// Generated by mcpp.tools.island for {0}. Do not edit.\n" + "//\n" + "// The island's boundary. Included by the device translation unit, which\n" + "// is compiled by a compiler mcpp did not resolve -- so the interface is\n" + "// `extern \"C\"`, because the two sides share no C++ ABI.\n" + "#ifndef {1}\n#define {1}\n" + "// Defined as nothing so the island can mark its entry points and\n" + "// still compile: the marker is for the generator to find, not for\n" + "// the compiler to act on.\n" + "#ifndef {2}\n#define {2}\n#endif\n" + "#ifdef __cplusplus\nextern \"C\" {{\n#endif\n\n", + by, guard, opt.marker); + for (auto const& e : entries) h += e + ";\n"; + h += "\n#ifdef __cplusplus\n}\n#endif\n#endif\n"; + + out.header_file = (dir / (opt.module_name + ".h")).string(); + out.include_dir = dir.string(); + if (!mcpp::plugins::surface::write_if_different(out.header_file, h)) { + std::cerr << std::format("mcpp.tools.island: cannot write {}\n", out.header_file); + return std::nullopt; + } + if (!opt.emit_module) return out; + + std::string m; + m += std::format("// Generated by mcpp.tools.island for {0}. Do not edit.\n" + "//\n" + "// The C++ side imports this instead of including the header. Names are\n" + "// re-exported rather than restated, so this file carries no second copy\n" + "// of a signature -- which at a C-linkage boundary is the copy that can\n" + "// disagree without anything noticing.\n" + "module;\n#include \"{1}\"\n" + "export module {2};\n\n", by, + std::filesystem::path(out.header_file).filename().string(), + opt.module_name); + for (auto const& n : names) m += std::format("export using ::{};\n", n); + + out.interface_file = (dir / (opt.module_name + ".cppm")).string(); + out.module_name = opt.module_name; + if (!mcpp::plugins::surface::write_if_different(out.interface_file, m)) { + std::cerr << std::format("mcpp.tools.island: cannot write {}\n", out.interface_file); + return std::nullopt; + } + return out; +} + +} // namespace mcpp::tools::island From c53013157fb7f65e8825b10665f33b55f22ceaf6 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 01:05:39 +0800 Subject: [PATCH 10/20] surface: the module root is the package's name, not its directory's `module_root_from_package` derived the module a consumer imports from the leaf of MCPP_MANIFEST_DIR, because that was the only thing a build program was told about itself. It is a directory name, and the two answers differ whenever a project lays a package out under a generic folder: mcpp's own `examples/09-heterogeneous/vulkan/app/` declares `name = "vulkan-saxpy"` and got `app.shaders`, which every `/app/` in a workspace would also claim. mcpp 2026.9.7.1 answers the right question, so this reads `mcpp::package_name()`. The directory leaf stays only for a value that cannot occur -- a manifest with no `[package] name` does not load -- so an empty answer would mean the contract moved. The fixture could not see the difference. Every project under tests/ had a package name equal to its directory name, so both derivations produced the same string. `spirv-module-consumer` now declares `name = "shader-app"` while keeping its directory, and its consumer imports `shader_app.shaders`; reverting this function makes that import unresolvable, which is the measurement the fixture was missing. --- .github/workflows/ci.yml | 19 ++++++++++++++----- src/plugins.cppm | 22 ++++++++++++++++++---- tests/spirv-module-consumer/mcpp.toml | 12 ++++++++++-- tests/spirv-module-consumer/src/main.cpp | 8 ++++++-- 4 files changed, 48 insertions(+), 13 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 07ff331..e483a93 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -123,14 +123,23 @@ jobs: fi echo "ok: the consumer names only the module" # And what it imports really is a generated module interface. - test -f target/.build-mcpp/out/spirv/spirv_module_consumer.shaders.cppm - grep -q '^export module spirv_module_consumer.shaders;' \ - target/.build-mcpp/out/spirv/spirv_module_consumer.shaders.cppm + # + # `shader_app`, while this directory is `spirv-module-consumer`. The + # module root is derived from `[package] name`, and the two are + # deliberately different here: a derivation that fell back to the + # manifest directory's leaf -- which is what this rule did before mcpp + # 2026.9.7.1 gave build programs `mcpp::package_name()` -- would write + # `spirv_module_consumer.shaders.cppm` and this test would find no file. + iface=target/.build-mcpp/out/spirv/shader_app.shaders.cppm + test -f "$iface" || { + echo "FAIL: no module interface at $iface" + find target/.build-mcpp/out/spirv -name '*.cppm' -print + exit 1; } + grep -q '^export module shader_app.shaders;' "$iface" # The interface must name no standard-library type: one that does # carries that header's templates in every consumer's BMI, measured at # 727 times the size of the std-free form. - if grep -n 'std::\|#include <' \ - target/.build-mcpp/out/spirv/spirv_module_consumer.shaders.cppm; then + if grep -n 'std::\|#include <' "$iface"; then echo "FAIL: the generated module interface names a standard-library type" exit 1 fi diff --git a/src/plugins.cppm b/src/plugins.cppm index 3bc2bf6..a51411c 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -628,11 +628,25 @@ inline kind default_surface() { // followed by the group's own segment, so two packages in one build cannot // claim the same module. inline std::string module_root_from_package() { + // THE PACKAGE'S NAME. NOT ITS DIRECTORY'S. + // + // These are different questions and they give different answers whenever a + // project lays a package out under a generic directory. The first version + // of this function asked the only question mcpp could answer -- the leaf of + // MCPP_MANIFEST_DIR -- so a package named `vulkan-saxpy` laid out as + // `vulkan/app/` generated `app.shaders`, and every `/app/` in a + // workspace claimed that same module. `mcpp::package_name()` was added to + // the build-program contract in mcpp 2026.9.7.1, which this package + // requires, so the right question is now askable. + // + // The directory leaf remains only as a value for the impossible case: a + // manifest without a `[package] name` does not load, so an empty answer + // here would mean the contract changed underneath. A rule that wants + // neither passes `options::module_name`, and this is not consulted. + std::string leaf{mcpp::package_name()}; + if (leaf.empty()) + leaf = std::filesystem::path(mcpp::manifest_dir()).filename().string(); std::string s; - // MCPP_MANIFEST_DIR's leaf is the package directory, which is the closest - // thing a build program is told about its own name. A rule that knows - // better passes `options::module_name` and this is not consulted. - const auto leaf = std::filesystem::path(mcpp::manifest_dir()).filename().string(); for (char c : leaf) s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; if (s.empty()) s = "app"; diff --git a/tests/spirv-module-consumer/mcpp.toml b/tests/spirv-module-consumer/mcpp.toml index c3fcdf7..41e0b55 100644 --- a/tests/spirv-module-consumer/mcpp.toml +++ b/tests/spirv-module-consumer/mcpp.toml @@ -10,11 +10,19 @@ # - the module name and the namespace are the same identifier path, derived # from the package with nothing declared for it. # +# THE PACKAGE NAME AND THE DIRECTORY NAME DIFFER ON PURPOSE. The rule derives +# the module root from `[package] name`, and it used to derive it from the leaf +# of the manifest directory -- two questions that give the same answer for every +# fixture whose name matches its folder, which was all of them. With the two +# separated, a derivation that fell back to the directory generates +# `spirv_module_consumer.shaders` and `src/main.cpp` fails to resolve its +# import. The directory keeps its name because the CI step addresses it by that. +# # No Vulkan runtime is involved, for the reason `spirv-consumer` records: the # property that identifies a SPIR-V module without executing it is the magic # number in its first word. [package] -name = "spirv-module-consumer" +name = "shader-app" namespace = "example" version = "0.1.0" description = "GLSL compute shaders reached through a generated C++ module" @@ -34,6 +42,6 @@ sources = [ { glob = "shaders/**/*.comp", accel = "vulkan1.2" }, ] -[targets.spirv-module-consumer] +[targets.shader-app] kind = "bin" main = "src/main.cpp" diff --git a/tests/spirv-module-consumer/src/main.cpp b/tests/spirv-module-consumer/src/main.cpp index 4897a5c..4a363ac 100644 --- a/tests/spirv-module-consumer/src/main.cpp +++ b/tests/spirv-module-consumer/src/main.cpp @@ -1,9 +1,13 @@ // Nothing here names a generated file. The shaders arrive through one import, // and the two that share a stem are told apart by the directory they came from. +// +// `shader_app`, not `spirv_module_consumer`: the module root comes from +// `[package] name`, and this project's directory is named differently on +// purpose so that the two derivations are distinguishable here. import std; -import spirv_module_consumer.shaders; +import shader_app.shaders; -namespace shaders = spirv_module_consumer::shaders; +namespace shaders = shader_app::shaders; namespace { From 38549b4a36114d20384dd70eb322e38f0dbf146b Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 01:17:56 +0800 Subject: [PATCH 11/20] tools-island: a seam has two halves, and scan is where they are compared The generator served one island. A seam has two implementations of one `extern "C"` boundary -- a device island and a host fallback -- and exactly one of them is in any link, which is the arrangement every example under `examples/09-heterogeneous` has. Handed only the device half, the generator produced an empty boundary in a `--no-accel` build and the C++ side failed on an unresolved name. `scan` now merges entries by name, so a build program hands it both files unconditionally: both exist on disk in either build, and which one is compiled is the manifest's decision rather than a condition the build program repeats. Two definitions of one name that declare it differently are refused there, naming both files and both signatures. Nothing else catches that. The two halves are never in one translation unit and never in one link, and C language linkage does not mangle, so a build with disagreeing halves is clean and the artifact reads its arguments by whichever signature it was compiled with. `scan` is the only point at which both texts exist at once. The fixture gains the host half under `cfg(not(accelerator = "vulkan"))`, so both legs are built and run, and a CI step perturbs one signature and asserts the refusal names both files. The README paragraph claiming the two could not disagree under `scan` was true only while `scan` took one file. --- .github/workflows/ci.yml | 81 ++++++++++++++++++++++++-- README.md | 43 +++++++++++--- tests/island-interface/build.mcpp | 10 ++++ tests/island-interface/mcpp.toml | 16 ++++- tests/island-interface/src/cpu/saxpy.c | 30 ++++++++++ tools/island.cppm | 44 +++++++++++++- 6 files changed, 208 insertions(+), 16 deletions(-) create mode 100644 tests/island-interface/src/cpu/saxpy.c diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index e483a93..aa7a9c0 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -243,16 +243,14 @@ jobs: # island reads the generated header through the compiler's # forced-include flag -- so a project using this generator has no # header in its source tree and no line naming one. - if grep -rn '#include' src/main.cpp src/kernels/saxpy.c; then + if grep -rn '#include' src/main.cpp src/kernels/saxpy.c src/cpu/saxpy.c; then echo "FAIL: a source names an include; the generator exists to remove it" exit 1 fi # AND THE CHECK THAT INCLUDE USED TO DO IS STILL THERE. The compiler - # sees the declarations, so a definition whose signature drifted fails - # where it was written. Verified by hand against this fixture with a - # hand-written entry list declaring `double n` where the definition - # says `unsigned n`: `saxpy.c:22: error: conflicting types for - # 'scale_device'`, at the definition rather than at the link. + # sees the declarations, so a definition whose signature drifted from + # the generated header fails where it was written rather than at the + # link. grep -q 'include' build.mcpp \ || { echo "FAIL: the build program no longer forces the header in"; exit 1; } # THE MARKER SELECTS, AND THE UNMARKED MUST NOT TRAVEL. The island has @@ -275,6 +273,77 @@ jobs: "$d/island_interface.kernels.h" echo "ok: one declaration, two generated artefacts, and only what was marked" + # THE SEAM'S OTHER HALF, AND THE DISAGREEMENT THAT NOTHING ELSE CATCHES. + # + # A seam has two implementations of one `extern "C"` boundary -- a device + # island and a host fallback -- and exactly one of them is in any link. + # `build.mcpp` hands `scan` both files unconditionally, because both exist + # on disk in either build and the manifest decides which is compiled. A + # build program that asked `mcpp::accel()` instead would carry a second + # copy of a decision the manifest already states. + # + # The CPU leg is not a variant of the step above. It is the build in which + # the device half is ABSENT, so a generator that read only the file it was + # given first would produce an empty boundary and the C++ side would fail + # on an unresolved name. + - name: the same boundary, generated from the half the build selected + working-directory: tests/island-interface + run: | + rm -rf target + "$MCPP" build --no-accel > cpu.log 2>&1 || { + echo "FAIL: the CPU leg does not build"; tail -25 cpu.log; exit 1; } + "$MCPP" run --no-accel | tee cpurun.log + grep -q '^out\[0\]=6 out\[1\]=12 out\[2\]=18 out\[3\]=24' cpurun.log + grep -q '^all ok' cpurun.log + # ONE set of declarations, not two. Both files mark the same two entry + # points; a merge that appended would declare each twice, which the + # module rejects as a redefinition and the header silently accepts. + d=target/.build-mcpp/out/island + n=$(grep -c '^export using ::' "$d/island_interface.kernels.cppm") + [ "$n" = 2 ] || { + echo "FAIL: the module re-exports $n names; the two halves declare 2" + cat "$d/island_interface.kernels.cppm"; exit 1; } + echo "ok: the CPU leg reaches the same generated boundary" + + # AND THE REVERSE LEG, WHICH IS THE POINT OF SCANNING BOTH. + # + # C language linkage does not mangle and the two halves are never in one + # link, so two declarations of one name that disagree produce a clean + # build and an artifact that reads its arguments by whichever signature it + # happened to be compiled with. `scan` is the only place both texts exist + # at once, so it is the only place that can refuse. + - name: two halves that disagree are refused where both texts exist + working-directory: tests/island-interface + run: | + cp src/cpu/saxpy.c /tmp/cpu_saxpy.bak + sed -i 's/^int scale_device(float a, float\* out, unsigned n) {/int scale_device(float a, float* out, double n) {/' \ + src/cpu/saxpy.c + grep -q 'double n' src/cpu/saxpy.c || { + echo "FAIL: the fixture was not perturbed; this step would assert nothing" + cp /tmp/cpu_saxpy.bak src/cpu/saxpy.c; exit 1; } + rm -rf target + set +e + "$MCPP" build > neg.log 2>&1 + rc=$? + set -e + cp /tmp/cpu_saxpy.bak src/cpu/saxpy.c + [ "$rc" != 0 ] || { echo "FAIL: the disagreement was accepted"; tail -20 neg.log; exit 1; } + grep -q 'declare it differently' neg.log || { + echo "FAIL: refused, but not for this reason"; tail -20 neg.log; exit 1; } + # Both file names, because a diagnostic naming one leaves the reader + # to find the other. + grep -q 'src/kernels/saxpy.c' neg.log + grep -q 'src/cpu/saxpy.c' neg.log + echo "ok: refused, naming both definitions and both signatures" + # And it builds again once they agree, so this step cannot leave a + # fixture that refuses everything. + rm -rf target + "$MCPP" build > restored.log 2>&1 || { + echo "FAIL: the fixture no longer builds after the perturbation was undone" + tail -20 restored.log; exit 1; } + rm -f neg.log restored.log cpu.log cpurun.log + echo "ok: and it builds again once the two agree" + # NO BUILD PROGRAM AT ALL, WHICH IS WHAT THE TWO KEYS BUY. # # `spirv-consumer` and `spirv-module-consumer` both write a `build.mcpp`, diff --git a/README.md b/README.md index 09f92e5..1c3ab34 100644 --- a/README.md +++ b/README.md @@ -274,8 +274,8 @@ The entry points are marked where they are defined, and the signature exists once: ```c -// src/kernels/saxpy.cu -#include "myapp.kernels.h" // generated; defines the marker as nothing +// src/kernels/saxpy.cu -- no include: the generated header arrives through the +// compiler's forced-include flag, which is what defines the marker as nothing. MCPP_EXPORT_C int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n) { ... } @@ -327,17 +327,44 @@ until mcpp writes it. **The check that include used to do is still there.** The compiler sees the declarations, so a definition whose signature drifted from its declaration fails -where it was written rather than at the link. Verified against the fixture with -a hand-written entry list declaring `double n` where the definition says -`unsigned n`: +where it was written rather than at the link: ``` src/kernels/saxpy.c:22:5: error: conflicting types for 'scale_device' ``` -That check is only reachable on the `emit` path, where the list and the -definition are separate things. Under `scan` the header is generated *from* the -definition, so the two cannot disagree at all. +**A seam has two halves, and `scan` is where they are compared.** A device +island and a host fallback implement one `extern "C"` boundary, and exactly one +of them is in any link. A build program hands `scan` both, unconditionally -- +both files exist on disk in either build, and which one is compiled is the +manifest's decision rather than a condition the build program repeats: + +```cpp +const std::vector islands{ + std::string(mcpp::manifest_dir()) + "/src/kernels/saxpy.cu", + std::string(mcpp::manifest_dir()) + "/src/cpu/saxpy.cpp", +}; +``` + +Entries are merged by name, so the two halves produce one set of declarations. +Two definitions of one name that declare it **differently** are refused, naming +both files and both signatures: + +``` +mcpp.tools.island: two definitions of `scale_device` declare it differently. + src/kernels/saxpy.c + int scale_device(float a, float* out, unsigned n) + src/cpu/saxpy.c + int scale_device(float a, float* out, double n) + C language linkage does not mangle, so these never meet at the link: + whichever one is in the artifact reads its arguments by its own signature. +``` + +Nothing else in the toolchain catches that. The two halves are never in one +translation unit and never in one link, and C linkage does not mangle, so a +build with disagreeing halves is clean and the artifact reads its arguments by +whichever signature it was compiled with. `scan` is the only point at which both +texts exist at once. **The declaration still exists once.** Without this, a project writes it twice -- in a header, and again wherever the C++ side reaches it. C language linkage does diff --git a/tests/island-interface/build.mcpp b/tests/island-interface/build.mcpp index b80ac47..1466d49 100644 --- a/tests/island-interface/build.mcpp +++ b/tests/island-interface/build.mcpp @@ -16,8 +16,14 @@ import mcpp.tools.island; // linkage does not mangle, so two copies that disagree are one symbol: the link // is clean and each side reads the arguments by its own ABI. int main() { + // BOTH HALVES, UNCONDITIONALLY. Which one the build compiles is the + // manifest's decision and this program does not repeat it: the declarations + // are the same either way, and `scan` merges them by entry name. A build + // program that asked `mcpp::accel()` here would be a second copy of a + // decision the manifest already states -- and the copy that goes stale. const std::vector islands{ std::string(mcpp::manifest_dir()) + "/src/kernels/saxpy.c", + std::string(mcpp::manifest_dir()) + "/src/cpu/saxpy.c", }; for (auto const& f : islands) mcpp::rerun_if_changed(f.c_str()); @@ -28,6 +34,10 @@ int main() { const auto entries = mcpp::tools::island::scan(islands, opt); if (!entries) return 1; + // TWO, not four. Both files mark the same two entry points, and a merge + // that appended instead would declare each one twice -- which the module + // rejects as a redefinition and the header does not, so the count is the + // criterion that catches it in both surfaces. if (entries->size() != 2) { std::cerr << std::format("expected two marked entry points, found {}\n", entries->size()); diff --git a/tests/island-interface/mcpp.toml b/tests/island-interface/mcpp.toml index 0f261c6..e7e5b6d 100644 --- a/tests/island-interface/mcpp.toml +++ b/tests/island-interface/mcpp.toml @@ -21,6 +21,7 @@ name = "island-interface" namespace = "example" version = "0.1.0" description = "An extern \"C\" island reached through a generated module" +accelerators = ["vulkan"] [language] standard = "c++23" @@ -30,8 +31,21 @@ import_std = true [build-dependencies.mcpp] plugins = { path = "../..", features = ["tools-island"], host-module = true } +# ONE BOUNDARY, TWO IMPLEMENTATIONS, AND EXACTLY ONE OF THEM IN ANY LINK. +# +# This is the arrangement every seam in `examples/09-heterogeneous` has, and it +# is what the generator has to serve without the project writing a condition: +# `build.mcpp` hands `scan` both files unconditionally, because both exist on +# disk in either build, and the manifest decides which one is compiled. [build] -sources = ["src/*.cpp", "src/kernels/*.c"] +accel = "vulkan1.2" +sources = ["src/*.cpp"] + +[target.'cfg(accelerator = "vulkan")'.build] +sources = ["src/kernels/*.c"] + +[target.'cfg(not(accelerator = "vulkan"))'.build] +sources = ["src/cpu/*.c"] [targets.island-interface] kind = "bin" diff --git a/tests/island-interface/src/cpu/saxpy.c b/tests/island-interface/src/cpu/saxpy.c new file mode 100644 index 0000000..13a07e8 --- /dev/null +++ b/tests/island-interface/src/cpu/saxpy.c @@ -0,0 +1,30 @@ +/* The host half of the same boundary. + * + * The signatures here and in `../kernels/saxpy.c` are the same declarations, + * and that is a property nothing else in the toolchain checks: C language + * linkage does not mangle, exactly one of these files is in any link, and two + * that disagreed would each read the arguments its own way. + * + * `mcpp::tools::island::scan` is handed both, so it is the one place where both + * texts exist at once -- and it refuses a disagreement there, with both file + * names, instead of leaving it to the run. + * + * No include here either. The generated boundary header arrives through the + * compiler's forced-include flag, the same way it reaches the device half. */ + +MCPP_EXPORT_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; +} + +MCPP_EXPORT_C +int scale_device(float a, float* out, unsigned n) { + for (unsigned i = 0; i < n; ++i) out[i] = a * out[i]; + return 0; +} + +/* The host half has its own internals, and they must not reach the boundary + * any more than the device half's do. */ +static int host_only_helper(int x) { return x - 1; } +int internal_device(int x) { return host_only_helper(x); } diff --git a/tools/island.cppm b/tools/island.cppm index 95e9ba1..5af24ca 100644 --- a/tools/island.cppm +++ b/tools/island.cppm @@ -159,9 +159,25 @@ inline std::string entry_name(std::string_view decl) { // point that produced no declaration would leave the island defining a function // nothing declares, and the consumer's failure would be an unresolved name in a // different file. +// +// SEVERAL SOURCES MAY DEFINE THE SAME ENTRY POINT, AND THEY MUST AGREE. +// +// That is the ordinary shape of a seam: a device island and a host fallback +// define one boundary, and exactly one of them is in any given link. A build +// program hands both to this function and gets one set of declarations back, +// so it does not have to ask which build it is in. +// +// Two definitions of one name whose declarations differ are REFUSED here, +// naming both files. Nothing else in the toolchain catches that: C language +// linkage does not mangle, so the two never meet at the link, and whichever +// one is present reads its arguments by its own idea of the signature. This is +// the only point at which both texts exist at once. inline std::optional> scan(std::span sources, const options& opt) { std::vector entries; + // Where each entry was found, for the disagreement diagnostic. Parallel to + // `entries`, which is the return value and cannot carry it. + std::vector origin; for (auto const& src : sources) { std::ifstream in(src); if (!in) { @@ -226,7 +242,33 @@ scan(std::span sources, const options& opt) { flat += c; } } - if (!flat.empty()) entries.push_back(std::move(flat)); + if (flat.empty()) continue; + + // MERGED BY ENTRY NAME. A second definition of a name already seen + // is either the same declaration -- the seam's two halves agreeing, + // which is the expected case -- or a disagreement that has to stop + // the build here. + const auto name = entry_name(flat); + std::size_t seen = entries.size(); + for (std::size_t k = 0; k < entries.size(); ++k) + if (entry_name(entries[k]) == name) { seen = k; break; } + + if (seen == entries.size()) { + entries.push_back(std::move(flat)); + origin.push_back(src); + continue; + } + if (entries[seen] == flat) continue; // both halves agree + + std::cerr << std::format( + "mcpp.tools.island: two definitions of `{}` declare it differently.\n" + " {}\n {}\n" + " {}\n {}\n" + " C language linkage does not mangle, so these never meet at the " + "link:\n whichever one is in the artifact reads its arguments by its " + "own signature.\n", + name, origin[seen], entries[seen], src, flat); + return std::nullopt; } } return entries; From bf81b356d4e8659fe65fc43af4f5f5215f363e13 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 01:20:11 +0800 Subject: [PATCH 12/20] tools-island: the fixture reaches the boundary through a seam module Every example under `examples/09-heterogeneous` puts a module of the project between the consumer and the island: `app.cppm` imports the boundary and turns pointers and a count back into spans, and it is the one place a backend can be exchanged. The fixture reached the generated module from `src/main.cpp`, which is not a module unit -- so its import said nothing about the case the examples are in, where one module interface written by the project imports a module interface written into the build directory during the same build. `src/app.cppm` is that seam, and `src/main.cpp` now names no boundary symbol at all. Both properties are asserted: a consumer that named `saxpy_device` would mean the seam is not hiding which island is underneath. Both legs build and run with it. --- .github/workflows/ci.yml | 13 ++++++++++- README.md | 7 ++++++ tests/island-interface/mcpp.toml | 2 +- tests/island-interface/src/app.cppm | 35 +++++++++++++++++++++++++++++ tests/island-interface/src/main.cpp | 34 ++++++++++++++++------------ 5 files changed, 75 insertions(+), 16 deletions(-) create mode 100644 tests/island-interface/src/app.cppm diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index aa7a9c0..a367a31 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -243,10 +243,21 @@ jobs: # island reads the generated header through the compiler's # forced-include flag -- so a project using this generator has no # header in its source tree and no line naming one. - if grep -rn '#include' src/main.cpp src/kernels/saxpy.c src/cpu/saxpy.c; then + if grep -rn '#include' src/main.cpp src/app.cppm src/kernels/saxpy.c src/cpu/saxpy.c; then echo "FAIL: a source names an include; the generator exists to remove it" exit 1 fi + # THE SEAM IS A MODULE OF THIS PROJECT IMPORTING THE GENERATED ONE. + # `main.cpp` is not a module unit, so its import said nothing about + # ordering two module interfaces where one is written during the + # build. `src/app.cppm` is, and it is the shape every example under + # examples/09-heterogeneous has. + grep -q '^import island_interface.kernels;' src/app.cppm \ + || { echo "FAIL: the seam no longer imports the generated module"; exit 1; } + if grep -q 'saxpy_device\|scale_device' src/main.cpp; then + echo "FAIL: the consumer names a boundary symbol; the seam exists to hide it" + exit 1 + fi # AND THE CHECK THAT INCLUDE USED TO DO IS STILL THERE. The compiler # sees the declarations, so a definition whose signature drifted from # the generated header fails where it was written rather than at the diff --git a/README.md b/README.md index 1c3ab34..3ea8382 100644 --- a/README.md +++ b/README.md @@ -297,6 +297,13 @@ Two files come out of that one marked declaration: the `extern "C"` header the device translation unit includes, guards and `__cplusplus` dance included, and the module the C++ side imports. +The C++ side is usually a **seam module** of the project rather than a consumer +directly: `app.cppm` imports the generated module and turns pointers and a count +back into spans, and it is the one place a backend can be exchanged. That means +one module interface of the project imports a module interface written into the +build directory during the same build; the ordering comes from the scan seeing +the import, and nothing has to be declared for it. + **Three layers, and each overrides the one above.** `scan` reads the marked declarations out of the island, which puts the signature beside the definition; `emit` takes a list directly, for entry points a scan cannot see; and a project diff --git a/tests/island-interface/mcpp.toml b/tests/island-interface/mcpp.toml index e7e5b6d..c41d526 100644 --- a/tests/island-interface/mcpp.toml +++ b/tests/island-interface/mcpp.toml @@ -39,7 +39,7 @@ plugins = { path = "../..", features = ["tools-island"], host-module = true } # disk in either build, and the manifest decides which one is compiled. [build] accel = "vulkan1.2" -sources = ["src/*.cpp"] +sources = ["src/*.cppm", "src/*.cpp"] [target.'cfg(accelerator = "vulkan")'.build] sources = ["src/kernels/*.c"] diff --git a/tests/island-interface/src/app.cppm b/tests/island-interface/src/app.cppm new file mode 100644 index 0000000..5052ac9 --- /dev/null +++ b/tests/island-interface/src/app.cppm @@ -0,0 +1,35 @@ +// The seam: a module of this project that imports the GENERATED module. +// +// This is the shape every example under `examples/09-heterogeneous` has, and it +// is the one the fixture was missing. `src/main.cpp` is not a module unit, so +// its `import` said nothing about the case that matters here -- a module +// interface written by this project importing a module interface written into +// the build directory during the same build. The two have to be ordered, and +// the ordering comes from the scan seeing the import rather than from anything +// this file declares. +// +// It is also where the boundary stops being C. Above this line callers pass +// spans; below it, pointers and a count, which is the one shape every device +// API agrees on. +export module island_interface.app; + +import std; +import island_interface.kernels; + +export namespace island_interface { + +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; +} + +bool scale(float a, std::span v) { + return scale_device(a, v.data(), static_cast(v.size())) == 0; +} + +} // namespace island_interface diff --git a/tests/island-interface/src/main.cpp b/tests/island-interface/src/main.cpp index 7170aea..ab85a28 100644 --- a/tests/island-interface/src/main.cpp +++ b/tests/island-interface/src/main.cpp @@ -1,28 +1,34 @@ -// The C++ side. It imports the generated module and includes nothing: the -// header exists for the island's compiler, which does not read modules. +// The consumer. It imports the SEAM, not the generated module, and includes +// nothing: the generated header exists for the island's compiler, which does +// not read modules. +// +// Nothing here names `saxpy_device`. That is the property the seam exists for: +// which island is underneath -- the device half or the host one -- is not +// visible from this file, and neither is the fact that a boundary was +// generated at all. import std; -import island_interface.kernels; +import island_interface.app; int main() { - float x[4] = {1, 2, 3, 4}; - float y[4] = {10, 20, 30, 40}; - float out[4] = {}; + const std::vector x{1, 2, 3, 4}; + const std::vector y{10, 20, 30, 40}; - if (saxpy_device(2.0f, x, y, out, 4) != 0) { - std::cout << "BAD: saxpy_device failed\n"; + auto out = island_interface::saxpy(2.0f, x, y); + if (!out) { + std::cout << "BAD: saxpy failed\n"; return 1; } - if (scale_device(0.5f, out, 4) != 0) { - std::cout << "BAD: scale_device failed\n"; + if (!island_interface::scale(0.5f, *out)) { + std::cout << "BAD: scale failed\n"; return 1; } // (2*1+10)/2, (2*2+20)/2, (2*3+30)/2, (2*4+40)/2 const float want[4] = {6, 12, 18, 24}; - bool ok = true; - for (int i = 0; i < 4; ++i) { - std::cout << std::format("out[{}]={} ", i, out[i]); - if (out[i] != want[i]) ok = false; + bool ok = out->size() == 4; + for (std::size_t i = 0; ok && i < out->size(); ++i) { + std::cout << std::format("out[{}]={} ", i, (*out)[i]); + if ((*out)[i] != want[i]) ok = false; } std::cout << "\n"; // Two entry points rather than one, because a generator that re-exported From 9a2ff4b1bb0a645d7f5c0201008b0bd6a54e4691 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 01:28:36 +0800 Subject: [PATCH 13/20] rules: cuda, hip and sycl gain the flag pass-through ascendc already had `mcpp.tools.island::force_include_flags` returns the flags that make a compiler read the generated boundary header before the island's first line, and three of the four device rules had no way to deliver them. Measured on `examples/09-heterogeneous/cuda`: the island is compiled by nvcc, or by clang in CUDA mode, on a command line the rule assembles -- a driver mcpp did not invoke, which inherits nothing from `mcpp::cflag` or `mcpp::cxxflag`. The project-wide channels are not merely too wide for the job, they are wrong for it. `cxxflag` forces the header into every C++ translation unit, including the seam `src/app.cppm`, and a module interface unit must begin with `export module` -- declarations ahead of that line are ill-formed. `mcpp.rules.ascendc` already had `options::flags`, appended last. The other three now carry the same field with the same position and the same pass-through-unexamined contract, because one rule with it and three without is the shape where a project's build program works for one backend and not the next -- which is what a seam exists to prevent. Verified end to end on this host: the cuda example's island compiles with `-include .../app.kernels.h` on its own command line, and the artifact runs on an RTX 4080. --- rules/cuda.cppm | 12 ++++++++++++ rules/hip.cppm | 12 ++++++++++++ rules/sycl.cppm | 12 ++++++++++++ 3 files changed, 36 insertions(+) diff --git a/rules/cuda.cppm b/rules/cuda.cppm index 27c4c44..964d3dc 100644 --- a/rules/cuda.cppm +++ b/rules/cuda.cppm @@ -97,6 +97,17 @@ struct options { // dependency's directory here, and it knows it only as the absolute path // `mcpp::dep_dir` answered with. std::vector includes; + // FLAGS FOR THE ISLAND'S COMPILER, PASSED THROUGH UNEXAMINED. + // + // A device compiler is a separate driver with its own command line, and + // `mcpp::cflag`/`mcpp::cxxflag` reach mcpp's compiler rather than this one. + // The case this exists for is `mcpp.tools.island`, whose + // `force_include_flags` makes the island read its generated boundary header + // before its first line -- so the island names no generated file and the + // project has no header of its own. Project-wide flags cannot do that job: + // forcing a header into every C++ translation unit puts declarations ahead + // of `export module`, which no module interface unit accepts. + std::vector flags; std::string out_dir = std::string(mcpp::out_dir()); }; @@ -728,6 +739,7 @@ inline std::vector plan(std::span sources, options opt for (auto const& inc : opt.includes) e.command.push_back("-I" + (std::filesystem::path(inc).is_absolute() ? inc : root + "/" + inc)); + for (auto const& f : opt.flags) e.command.push_back(f); e.command.insert(e.command.end(), { "-c", root + "/" + src, "-o", obj }); e.inputs = { root + "/" + src }; e.outputs = { obj }; diff --git a/rules/hip.cppm b/rules/hip.cppm index eb5a8dd..f216889 100644 --- a/rules/hip.cppm +++ b/rules/hip.cppm @@ -69,6 +69,17 @@ struct options { // the form `mcpp::dep_dir` answers with -- a device compiler is a separate // driver and inherits nothing from the C++ side's include configuration. std::vector includes; + // FLAGS FOR THE ISLAND'S COMPILER, PASSED THROUGH UNEXAMINED. + // + // A device compiler is a separate driver with its own command line, and + // `mcpp::cflag`/`mcpp::cxxflag` reach mcpp's compiler rather than this one. + // The case this exists for is `mcpp.tools.island`, whose + // `force_include_flags` makes the island read its generated boundary header + // before its first line -- so the island names no generated file and the + // project has no header of its own. Project-wide flags cannot do that job: + // forcing a header into every C++ translation unit puts declarations ahead + // of `export module`, which no module interface unit accepts. + std::vector flags; std::string out_dir = std::string(mcpp::out_dir()); }; @@ -420,6 +431,7 @@ inline std::vector plan(std::span sources, options opt for (auto const& inc : opt.includes) e.command.push_back("-I" + (std::filesystem::path(inc).is_absolute() ? inc : root + "/" + inc)); + for (auto const& f : opt.flags) e.command.push_back(f); e.command.insert(e.command.end(), { "-c", root + "/" + src, "-o", obj }); e.inputs = { root + "/" + src }; e.outputs = { obj }; diff --git a/rules/sycl.cppm b/rules/sycl.cppm index cd204d3..fe0ba89 100644 --- a/rules/sycl.cppm +++ b/rules/sycl.cppm @@ -98,6 +98,17 @@ struct options { // Header search paths for the island. Relative entries resolve against the // package root; an ABSOLUTE entry is passed through unchanged. std::vector includes; + // FLAGS FOR THE ISLAND'S COMPILER, PASSED THROUGH UNEXAMINED. + // + // A device compiler is a separate driver with its own command line, and + // `mcpp::cflag`/`mcpp::cxxflag` reach mcpp's compiler rather than this one. + // The case this exists for is `mcpp.tools.island`, whose + // `force_include_flags` makes the island read its generated boundary header + // before its first line -- so the island names no generated file and the + // project has no header of its own. Project-wide flags cannot do that job: + // forcing a header into every C++ translation unit puts declarations ahead + // of `export module`, which no module interface unit accepts. + std::vector flags; // An explicit compiler path wins over the payload. Set it when a project // pins a DPC++ other than the one the workspace installed. std::string compiler; @@ -483,6 +494,7 @@ inline std::vector plan(std::span sources, options opt for (auto const& inc : opt.includes) e.command.push_back("-I" + (std::filesystem::path(inc).is_absolute() ? inc : root + "/" + inc)); + for (auto const& f : opt.flags) e.command.push_back(f); // `-x c++` is not optional. `.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 From f20bb111876abb683763f642834cf1599b666628 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 01:58:05 +0800 Subject: [PATCH 14/20] README: two claims left over from the design that was abandoned The `rules-slang` row said its engine need is `.slang` in the device-source table. That was the first fix and it inverts the dependency; what shipped is the feature declaring `device_extensions` and `rule_module` itself, with `.slang` REMOVED from the engine's table. Its absence from the list of per-member floors is now stated rather than left as a gap, because that absence is the criterion for the whole arrangement. The shared floor moves to 2026.9.7.1, and the two usage snippets pin 0.3.0. --- README.md | 23 ++++++++++++++++++----- 1 file changed, 18 insertions(+), 5 deletions(-) diff --git a/README.md b/README.md index 3ea8382..fe71f9c 100644 --- a/README.md +++ b/README.md @@ -6,7 +6,7 @@ imports each one from `build.mcpp` under the module name the member declares. ```toml [build-dependencies.mcpp] -plugins = { version = "0.2.3", features = ["rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-spirv"], host-module = true } ``` `[build-dependencies]`, not `[dependencies]`. The two keys answer separate @@ -49,7 +49,7 @@ engine's own module family and is not used here. | `rules-ascendc` | `mcpp.rules.ascendc` | 2026.9.6.6 | `[build] accel = "ascend8.5+{dav-c220}"`, a constrained glob for `*.asc`. Compiles with BiSheng in MIXED mode, so the object carries the device binary and a host-callable launcher and joins the ordinary link -- no registration file and no device-link step. Its own engine needs are `.asc` in the device-source table and `mcpp::link_flag` for the `-rpath-link` the toolkit's shared libraries require, both 2026.9.6.5 | | `rules-cuda` | `mcpp.rules.cuda` | 2026.9.6.6 | `[build] accel = "cuda…"`, a constrained glob for `*.cu`; the clang route with an LLVM toolchain, the nvcc route with a GCC one | | `rules-hip` | `mcpp.rules.hip` | 2026.9.6.6 | `[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-slang` | `mcpp.rules.slang` | see below | `[build] accel = "vulkan1.2"`, a constrained glob for `*.slang`. Slang is a different language from GLSL rather than a second driver for it -- its own module system, generics, and targets beyond SPIR-V -- so it is a rule of its own. Its engine need is `.slang` in the device-source table | +| `rules-slang` | `mcpp.rules.slang` | 2026.9.7.1 | `[build] accel = "vulkan1.2"`, a constrained glob for `*.slang`. Slang is a different language from GLSL rather than a second driver for it -- its own module system, generics, and targets beyond SPIR-V -- so it is a rule of its own. `.slang` is **not** in the engine's device-source table: this feature declares `device_extensions = [".slang"]` and `rule_module = "mcpp.rules.slang"`, and the engine routes it from there. That is the criterion for the whole arrangement -- a new device language costs no engine release | | `rules-spirv` | `mcpp.rules.spirv` | 2026.9.6.6 | `[build] accel = "vulkan1.2"`, a constrained glob for the shader stages; compiles each shader through a `role = "source"` action and states which of the two compilers produced it | | `rules-sycl` | `mcpp.rules.sycl` | 2026.9.6.6 | `[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. Its own engine need is `.sycl` in the device-source table, 2026.9.6.1 | | `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 | @@ -61,7 +61,7 @@ A project names the rule and nothing else: ```toml [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-cuda"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-cuda"], host-module = true } ``` The payloads each rule drives are declared **here**, under the feature that @@ -127,8 +127,15 @@ device source that reached no action, naming the file. That is the engine's half of this rule and it needs 2026.9.6.5. The floor is the mcpp release whose engine carries what the member relies on. -From 0.2.4 every rule shares one: **2026.9.6.6**, the release in which a payload -a DEPENDENCY declared is both installed and answerable. Before it a rule could +From 0.3.0 every rule shares one: **2026.9.7.1**, the release that reads +`device_extensions` and `rule_module`, reports `[language] modules` and the +package's own name to a build program, and writes the build program a declared +rule set describes. A client below it does not get a degraded surface; it gets a +build in which the rules never route -- the file falls through to the ordinary +source scan and mcpp says it has no role for the extension. + +The previous shared floor was 2026.9.6.6, the release in which a payload a +DEPENDENCY declared is both installed and answerable. Before it a rule could declare `>=8.5.0`, have it installed, and still be told by `xpkg_dir` that nothing was there -- which is why each rule's list used to be repeated in every project that used it. The earlier per-member floors are still the floors of the @@ -137,6 +144,12 @@ table, 2026.9.5.3; `tools-embed` needs the fast path to compare a declared file input, 2026.9.5.4; `rules-sycl` needs `.sycl` in that table, 2026.9.6.1), and they are all below the shared one. +`rules-slang` is the member that does NOT appear in that list, and its absence +is the point: `.slang` is in no engine table at any version. The feature +declares the extension and the module that compiles it, so the release it needs +is the one that reads those two keys rather than the one that would have carried +its extension. + 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. From 6740b6e506c41aeee5ab8d4d92f37624b9572067 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 02:02:49 +0800 Subject: [PATCH 15/20] surface: a generated name may not be a C++ keyword Six sites across four files ran the same character filter -- letters, digits and `_` survive, everything else becomes `_`, a leading digit gets a `_` -- and none of them asked whether the result is reserved. `default`, `template`, `operator`, `private` and `union` are valid identifiers to that filter, and `shaders/default/` is an ordinary name for a shader directory. Measured: renaming the fixture's `shaders/a` to `shaders/default` makes the generator write namespace default { and the build fails at `expected identifier before 'default'`, inside a generated file, on a line the author of that directory has never opened. `surface::identifier` does all three transformations in one place, and a reserved result gets a TRAILING underscore -- a leading one is itself reserved at namespace scope, so prefixing would trade one reserved name for another. Namespace segments in `rules-spirv` and `rules-slang` go through it, as does `tools-embed`'s accessor name. The accessor is sanitised where it is EMITTED rather than in each producer. A rule that builds `item::identifier` from a file stem cannot know it has produced `my-shader_comp` or `default` until it reaches the one line that writes the function's name. `accessor_base` and the `_spv` symbols are left alone: both carry fixed affixes and cannot come out reserved, and routing them would rename symbols already published. The fixture could not see any of this: `a` and `b` are ordinary identifiers, so both filters produce the same file. It gains `shaders/default/`, and CI asserts the generated interface contains `namespace default_ {`. --- .github/workflows/ci.yml | 10 +++ rules/slang.cppm | 9 +-- rules/spirv.cppm | 9 +-- src/plugins.cppm | 75 +++++++++++++++++-- tests/spirv-module-consumer/mcpp.toml | 6 +- .../shaders/default/scale.comp | 15 ++++ tests/spirv-module-consumer/src/main.cpp | 8 ++ tools/embed.cppm | 10 +-- 8 files changed, 118 insertions(+), 24 deletions(-) create mode 100644 tests/spirv-module-consumer/shaders/default/scale.comp diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index a367a31..64c8661 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -115,7 +115,17 @@ jobs: "$MCPP" run | tee run.log grep -q '^a/scale.comp: magic=07230203' run.log grep -q '^b/scale.comp: magic=07230203' run.log + grep -q '^default/scale.comp: magic=07230203' run.log grep -q '^all ok' run.log + # A DIRECTORY NAMED AFTER A C++ KEYWORD. `a` and `b` are ordinary + # identifiers and cannot tell a name filter that knows the keywords + # from one that does not; without the guard this is + # `namespace default {` and the generated file does not parse. + grep -q '^namespace default_ {' \ + target/.build-mcpp/out/spirv/shader_app.shaders.cppm \ + || { echo "FAIL: the keyword directory did not get its trailing underscore" + grep -n namespace target/.build-mcpp/out/spirv/shader_app.shaders.cppm + exit 1; } # The point of the surface: a consumer names no generated file. if grep -rn 'scale_comp\.h\|\.inc"' src/; then echo "FAIL: a consumer source names a generated file" diff --git a/rules/slang.cppm b/rules/slang.cppm index 7c9bad6..932be07 100644 --- a/rules/slang.cppm +++ b/rules/slang.cppm @@ -267,11 +267,10 @@ inline std::vector namespace_of(std::string_view src, std::string_v for (auto const& part : std::filesystem::path(dir)) { auto s = part.string(); if (s.empty() || s == "." || s == "/") continue; - std::string seg; - for (char c : s) - seg += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; - if (std::isdigit(static_cast(seg.front()))) seg.insert(seg.begin(), '_'); - out.push_back(std::move(seg)); + // Through the lib root, which is the one place that knows a segment + // may not be a keyword: `shaders/default/` is an ordinary directory + // name and `namespace default {` is not a namespace. + out.push_back(mcpp::plugins::surface::identifier(s, "dir")); } return out; } diff --git a/rules/spirv.cppm b/rules/spirv.cppm index 2bdeda5..4880587 100644 --- a/rules/spirv.cppm +++ b/rules/spirv.cppm @@ -477,11 +477,10 @@ inline std::vector namespace_of(std::string_view src, std::string_v for (auto const& part : std::filesystem::path(dir)) { auto s = part.string(); if (s.empty() || s == "." || s == "/") continue; - std::string seg; - for (char c : s) - seg += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; - if (std::isdigit(static_cast(seg.front()))) seg.insert(seg.begin(), '_'); - out.push_back(std::move(seg)); + // Through the lib root, which is the one place that knows a segment + // may not be a keyword: `shaders/default/` is an ordinary directory + // name and `namespace default {` is not a namespace. + out.push_back(mcpp::plugins::surface::identifier(s, "dir")); } return out; } diff --git a/src/plugins.cppm b/src/plugins.cppm index a51411c..2ddca20 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -267,6 +267,68 @@ inline std::string accessor_base(const options& opt, const item& it) { return s; } +// A GENERATED NAME THE C++ COMPILER WILL ACCEPT. +// +// Three transformations, and the third is the one every hand-rolled copy of +// this function was missing. Non-identifier characters become `_`; a leading +// digit gets a `_` in front; and a result that is a KEYWORD gets a trailing `_`. +// +// The keyword case is not hypothetical. The first two rules accept `default`, +// `template`, `operator`, `private` and `union` unchanged -- they are valid +// identifiers to a character filter and reserved to the compiler -- and +// `shaders/default/` is an ordinary name for a shader directory. What it +// produced was `namespace default {` in a generated file, and an error naming a +// line its author never wrote. +// +// TRAILING `_`, not a prefix: `_default` is reserved at namespace scope +// (a leading underscore in the global namespace), and prefixing would trade one +// reserved name for another. +// +// The list is the keywords of the standard this collection targets. A word that +// is contextual rather than reserved (`final`, `override`, `import`, `module`) +// is a legal identifier and is left alone. +inline bool is_cxx_keyword(std::string_view w) { + static constexpr std::string_view kWords[] = { + "alignas", "alignof", "and", "and_eq", "asm", "auto", "bitand", "bitor", + "bool", "break", "case", "catch", "char", "char8_t", "char16_t", + "char32_t", "class", "compl", "concept", "const", "consteval", + "constexpr", "constinit", "const_cast", "continue", "co_await", + "co_return", "co_yield", "decltype", "default", "delete", "do", "double", + "dynamic_cast", "else", "enum", "explicit", "export", "extern", "false", + "float", "for", "friend", "goto", "if", "inline", "int", "long", + "mutable", "namespace", "new", "noexcept", "not", "not_eq", "nullptr", + "operator", "or", "or_eq", "private", "protected", "public", "register", + "reinterpret_cast", "requires", "return", "short", "signed", "sizeof", + "static", "static_assert", "static_cast", "struct", "switch", + "template", "this", "thread_local", "throw", "true", "try", "typedef", + "typeid", "typename", "union", "unsigned", "using", "virtual", "void", + "volatile", "wchar_t", "while", "xor", "xor_eq", + }; + for (auto k : kWords) if (k == w) return true; + return false; +} + +// `fallback` is used when the input sanitises to nothing, which a file named +// only in punctuation does. +// +// INDEXED RATHER THAN A RANGE-FOR over the string, for the reason recorded in +// `mcpp.tools.island`: iterating a `std::string` inside an exported inline +// function makes GCC 16 instantiate its iterator in this BMI, and a consumer's +// build program then fails to compile on `always_inline` in a header naming +// neither this file nor this loop. +inline std::string identifier(std::string_view raw, std::string_view fallback) { + std::string s; + for (std::size_t i = 0; i < raw.size(); ++i) { + const char c = raw[i]; + s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; + } + if (s.empty()) s = std::string(fallback); + if (!s.empty() && std::isdigit(static_cast(s.front()))) + s.insert(s.begin(), '_'); + if (is_cxx_keyword(s)) s += '_'; + return s; +} + inline const char* element_type(element e) { return e == element::word32 ? "unsigned int" : "unsigned char"; } @@ -354,8 +416,12 @@ inline std::string declarations(std::span items, const options& opt) for (auto const& it : items) { reopen(it.name_space); const auto base = accessor_base(opt, it); + // THROUGH `identifier`, HERE RATHER THAN IN EACH PRODUCER. This is the + // one line that turns `item::identifier` into something a compiler + // parses, and a rule that builds the field from a file stem cannot know + // it has produced `my-shader_comp` or `default` until it gets here. s += std::format("inline payload {}() {{ return {{ {}_data(), {}_size() }}; }}\n", - it.identifier, base, base); + identifier(it.identifier, "payload"), base, base); } reopen({}); return s; @@ -646,12 +712,7 @@ inline std::string module_root_from_package() { std::string leaf{mcpp::package_name()}; if (leaf.empty()) leaf = std::filesystem::path(mcpp::manifest_dir()).filename().string(); - std::string s; - for (char c : leaf) - s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; - if (s.empty()) s = "app"; - if (std::isdigit(static_cast(s.front()))) s.insert(s.begin(), '_'); - return s; + return identifier(leaf, "app"); } } // namespace mcpp::plugins::surface diff --git a/tests/spirv-module-consumer/mcpp.toml b/tests/spirv-module-consumer/mcpp.toml index 41e0b55..1101b8c 100644 --- a/tests/spirv-module-consumer/mcpp.toml +++ b/tests/spirv-module-consumer/mcpp.toml @@ -6,7 +6,11 @@ # header anywhere in its own sources; # - two shaders with one stem in two directories are no longer a collision: # `shaders/a/scale.comp` and `shaders/b/scale.comp` land in -# `spirv_module_consumer::shaders::a` and `::b`; +# `shader_app::shaders::a` and `::b`; +# - a directory whose name is a C++ KEYWORD is reached as `default_`. `a` and +# `b` are ordinary identifiers and cannot tell a name filter that knows the +# keywords from one that does not; `shaders/default/` can, and without the +# guard the generator writes `namespace default {`; # - the module name and the namespace are the same identifier path, derived # from the package with nothing declared for it. # diff --git a/tests/spirv-module-consumer/shaders/default/scale.comp b/tests/spirv-module-consumer/shaders/default/scale.comp new file mode 100644 index 0000000..360eb26 --- /dev/null +++ b/tests/spirv-module-consumer/shaders/default/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-module-consumer/src/main.cpp b/tests/spirv-module-consumer/src/main.cpp index 4a363ac..995d418 100644 --- a/tests/spirv-module-consumer/src/main.cpp +++ b/tests/spirv-module-consumer/src/main.cpp @@ -28,6 +28,14 @@ int main() { bool ok = true; ok &= check("a/scale.comp", shaders::a::scale_comp()); ok &= check("b/scale.comp", shaders::b::scale_comp()); + // A DIRECTORY NAMED AFTER A C++ KEYWORD, reached as `default_`. + // + // `a` and `b` are ordinary identifiers, so neither could tell a name filter + // that knows the keywords from one that does not. `shaders/default/` is an + // ordinary name for a shader directory and `namespace default {` is not a + // namespace: without the trailing underscore the generator writes a file + // that does not parse, and the error names a line nobody wrote. + ok &= check("default/scale.comp", shaders::default_::scale_comp()); // The two shaders are the same source in two directories, so they must // produce identical modules -- and they must be two distinct objects, not diff --git a/tools/embed.cppm b/tools/embed.cppm index eb562d8..4e7a85a 100644 --- a/tools/embed.cppm +++ b/tools/embed.cppm @@ -89,13 +89,11 @@ struct options { // ---- internals ------------------------------------------------------------- +// The accessor's own name, so a file called `default.bin` must not produce +// `default()`. The lib root owns that decision; this is the one caller that +// needs it here. inline std::string sanitise(std::string_view stem) { - std::string s; - for (char c : stem) - s += (std::isalnum(static_cast(c)) || c == '_') ? c : '_'; - if (s.empty()) s = "data"; - if (std::isdigit(static_cast(s.front()))) s.insert(s.begin(), '_'); - return s; + return mcpp::plugins::surface::identifier(stem, "data"); } inline std::string default_dir() { From b3d665c7eb38d5d06f04d28da1f21c45fa6ff436 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 02:04:17 +0800 Subject: [PATCH 16/20] surface: a module name whose segment is not an identifier is refused The module name is the one part of this surface a project writes itself, and each of its segments becomes a namespace. `app.my-shaders` produced `namespace my-shaders {` in a generated file, so the error named a line nobody wrote and a file the author had never opened. Refused where it was written, naming the offending segment and the spelling that would work: mcpp.plugins.surface: `app.my-shaders` is not a usable module name: the segment `my-shaders` is not a C++ identifier. Each segment becomes a namespace, so it has to be one; `my_shaders` would work. The predicate is `identifier()` applied to the segment and compared with it, so the refusal and the derivation cannot disagree about what an identifier is -- including the keyword case. --- src/plugins.cppm | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/src/plugins.cppm b/src/plugins.cppm index 2ddca20..344df85 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -503,6 +503,20 @@ inline std::optional emit(std::span items, const options& o } const auto segs = split_module_name(opt.module_name); + // REFUSED HERE, NOT IN THE GENERATED FILE. The module name is the one part + // of this surface a project writes itself, and each of its segments becomes + // a namespace. A segment that is not an identifier -- empty, starting with + // a digit, carrying a `-`, or reserved -- produces a generated file that + // does not parse, and the error then names a line nobody wrote. + for (std::size_t i = 0; i < segs.size(); ++i) { + if (segs[i] == identifier(segs[i], "")) continue; + std::cerr << std::format( + "mcpp.plugins.surface: `{}` is not a usable module name: the segment `{}` " + "is not a C++ identifier.\n Each segment becomes a namespace, so it has " + "to be one; `{}` would work.\n", + opt.module_name, segs[i], identifier(segs[i], "part")); + return std::nullopt; + } const auto by = opt.produced_by.empty() ? std::string("mcpp.plugins.surface") : opt.produced_by; From 0f491bb3eebf6eae70a2daf786cb84e8ded7e932 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 02:13:33 +0800 Subject: [PATCH 17/20] ci: the denominator reads feature names off the section headers `all-rules-compile` asserts EVERY feature is compiled for this host, and "every" is a list the fixture carries -- so the list is compared against the package's own features before the build. That extractor was written for [features] rules-cuda = { sources = ["rules/cuda.cppm"] } and 0.3.0 spells them as sub-tables, because a feature now carries `device_extensions` and `rule_module` as well and TOML 1.0 forbids a multi-line inline table. `grep -oE '^[a-z-]+ +='` then matched the keys INSIDE those tables: package: sources sources sources sources fixture: rules-ascendc rules-cuda rules-hip rules-slang ... which reads as "the fixture is wrong" while the fixture is complete and the check is what broke. A denominator that cannot state the set it measures against is worse than none: it fails loudly for the wrong reason, and the next person edits the fixture. The names come off `[features.]` headers now, and the empty-result guard says the manifest's spelling changed rather than that the package has no features. Both copies of the check are replaced; fixing one would leave the other failing on the next host. --- .github/workflows/ci.yml | 38 ++++++++++++++++++++++++++++++++------ 1 file changed, 32 insertions(+), 6 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 64c8661..630da16 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -566,11 +566,24 @@ jobs: # added to this package would be covered by a step whose name says it # already is. The list is therefore compared against the package's # own `[features]` before the build, and the two must be equal. - feats=$(sed -n '/^\[features\]/,/^# /p' mcpp.toml \ - | grep -oE '^[a-z-]+ +=' | sed 's/ *=//' | grep -v '^default$' | sort) + # OFF THE SECTION HEADERS, which is where a feature's name is + # since 0.3.0: a feature carries `device_extensions` and + # `rule_module` as well as `sources`, and TOML 1.0 forbids a + # multi-line inline table, so `[features.]` replaced + # ` = { ... }`. Reading keys out of the table BODY matched + # `sources` once per feature and reported the fixture as + # incomplete while the extractor was what broke. + feats=$(grep -oE '^\[features\.[a-z0-9-]+\]' mcpp.toml \ + | sed 's/^\[features\.//; s/\]$//' | sort) used=$(sed -n '/features = \[/,/\], host-module/p' tests/all-rules-compile/mcpp.toml \ | grep -oE '"[a-z-]+"' | tr -d '"' | sort) - [ -n "$feats" ] || { echo "FAIL: read no features out of mcpp.toml"; exit 1; } + [ -n "$feats" ] || { + echo "FAIL: read no features out of mcpp.toml. The extractor above" + echo " expects [features.] sections and found none, so the" + echo " manifest spelling changed -- the package has not stopped" + echo " having features." + grep -n '^\[features' mcpp.toml | head + exit 1; } [ "$feats" = "$used" ] || { echo "FAIL: the fixture does not name every published feature" echo " package: $(echo $feats)" @@ -850,11 +863,24 @@ jobs: # added to this package would be covered by a step whose name says it # already is. The list is therefore compared against the package's # own `[features]` before the build, and the two must be equal. - feats=$(sed -n '/^\[features\]/,/^# /p' mcpp.toml \ - | grep -oE '^[a-z-]+ +=' | sed 's/ *=//' | grep -v '^default$' | sort) + # OFF THE SECTION HEADERS, which is where a feature's name is + # since 0.3.0: a feature carries `device_extensions` and + # `rule_module` as well as `sources`, and TOML 1.0 forbids a + # multi-line inline table, so `[features.]` replaced + # ` = { ... }`. Reading keys out of the table BODY matched + # `sources` once per feature and reported the fixture as + # incomplete while the extractor was what broke. + feats=$(grep -oE '^\[features\.[a-z0-9-]+\]' mcpp.toml \ + | sed 's/^\[features\.//; s/\]$//' | sort) used=$(sed -n '/features = \[/,/\], host-module/p' tests/all-rules-compile/mcpp.toml \ | grep -oE '"[a-z-]+"' | tr -d '"' | sort) - [ -n "$feats" ] || { echo "FAIL: read no features out of mcpp.toml"; exit 1; } + [ -n "$feats" ] || { + echo "FAIL: read no features out of mcpp.toml. The extractor above" + echo " expects [features.] sections and found none, so the" + echo " manifest spelling changed -- the package has not stopped" + echo " having features." + grep -n '^\[features' mcpp.toml | head + exit 1; } [ "$feats" = "$used" ] || { echo "FAIL: the fixture does not name every published feature" echo " package: $(echo $feats)" From 052aa44cba99a1893e5b73a4695ed2d7ff96b8dc Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 02:19:08 +0800 Subject: [PATCH 18/20] ci: the duplicate-stem step asserts what replaced the refusal Two shaders differing only by directory used to produce one output -- the name was the stem and the stage, so `a/x.vert` and `b/x.vert` both gave `x_vert.h` declaring `x_vert_spv` -- and the rule refused. The shader's directory is now part of the generated file's path and of the array's name, so there is nothing to refuse and the step failed with FAIL: two shaders mapping to one output were accepted which is the correct new behaviour reported as a defect. That change was not made for ergonomics: two byte-identical generated headers collapse under GCC's `#pragma once`, which treats two files of the same size and content as the same file, so the second include silently did nothing and both accessors returned one array. The step now asserts the property that made the refusal unnecessary -- `scale_comp_spv` and `dup_scale_comp_spv` are both declared -- and the criterion is the generated names rather than the run, because a build that collapsed them would also produce a working binary. The refusal stays in the rule. Its condition is now two shaders with one name in one directory, which a single glob cannot construct and an explicit `options::base_dir` can. Verified locally against the released 2026.9.7.1 binary. --- .github/workflows/ci.yml | 58 +++++++++++++++++++++++----------------- 1 file changed, 34 insertions(+), 24 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 630da16..931fe95 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -176,15 +176,23 @@ jobs: # TWO SHADERS THAT DIFFER ONLY BY DIRECTORY. # - # The output name is the stem and the stage, so `a/x.vert` and `b/x.vert` - # both produce `x_vert.h` declaring `x_vert_spv`. Before this rule checked - # for it, ninja caught the duplicate output -- so it was never silent, but - # the message named the generated file and neither shader, arrived as a - # graph-loading failure rather than as this rule's refusal, and stated no - # way out. A project with one shader per stage never meets it, which is - # how it survived; a graphics project organising shaders by purpose is the - # first to have two. - - name: two shaders with one stem are refused, naming both + # They used to produce one output -- the name was the stem and the stage, + # so `a/x.vert` and `b/x.vert` both gave `x_vert.h` declaring + # `x_vert_spv` -- and the rule refused, because ninja's duplicate-output + # error named the generated file and neither shader. + # + # THE COLLISION IS GONE, and this step asserts what replaced it. The + # shader's directory is part of the generated file's path and of the + # array's name, so the two are two things. That change was not made for + # ergonomics: two byte-identical generated headers collapse under GCC's + # `#pragma once`, which treats two files of the same size and content as + # the same file, so the second include silently did nothing and both + # accessors returned one array. + # + # The rule's refusal stays, and its condition is now two shaders with one + # name in ONE directory. A single glob cannot construct that; an explicit + # `options::base_dir` can. + - name: two shaders with one stem are two payloads, not one working-directory: tests/spirv-consumer run: | set -e @@ -194,21 +202,23 @@ jobs: cp mcpp.toml mcpp.toml.bak sed -i 's#shaders/\*.comp#shaders/**/*.comp#' mcpp.toml rm -rf target - if "$MCPP" build > dup.log 2>&1; then - echo "FAIL: two shaders mapping to one output were accepted" - exit 1 - fi - for needle in 'two shaders map to one output' 'shaders/scale.comp' \ - 'shaders/dup/scale.comp' 'scale_comp_spv' 'rename one of them'; do - grep -q -- "$needle" dup.log || { - echo "FAIL: the refusal does not mention '$needle'" - tail -20 dup.log - exit 1 - } - done - echo "ok: refused, naming both shaders, the symbol and the way out" - # …and the fixture still builds once the duplicate is gone, so this - # step cannot leave a project that refuses everything. + "$MCPP" build > dup.log 2>&1 || { + echo "FAIL: two shaders differing only by directory were refused" + tail -20 dup.log; mv mcpp.toml.bak mcpp.toml; rm -rf shaders/dup; exit 1; } + # TWO declarations, and the directory is in the second one's name. + # A build that collapsed them would also produce a working binary, so + # the criterion is the generated names rather than the run. + d=target/.build-mcpp/out/spirv + find "$d" -name '*.h' | sort | sed 's/^/ /' + grep -rq 'scale_comp_spv' "$d" || { + echo "FAIL: the top-level shader lost its declaration"; exit 1; } + grep -rq 'dup_scale_comp_spv' "$d" || { + echo "FAIL: the shader in dup/ did not take its directory into its name" + grep -rho '[a-z_]*scale_comp_spv' "$d" | sort -u + exit 1; } + echo "ok: two directories sharing a stem are two payloads" + # …and the fixture still builds with one shader per stem, so this step + # cannot leave a project that refuses everything. mv mcpp.toml.bak mcpp.toml rm -rf shaders/dup target "$MCPP" build > restored.log 2>&1 || { From 7eb610d1cf5cc639c3d01b09a06decd3f371ec55 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 02:21:07 +0800 Subject: [PATCH 19/20] ci: the consumer check asks about code, not about text The step asserts that `tests/island-interface/src/main.cpp` names no boundary symbol, because the seam exists to hide which island is underneath. It was a plain grep, and that file EXPLAINS the property in a comment -- "Nothing here names \`saxpy_device\`" -- so the criterion reported the file for saying what it does. Comment lines are stripped first. The difference is between asking whether the file mentions the name and whether it uses it, and only the second is the property. Measured both ways: `src/main.cpp` has zero uses after stripping and `src/app.cppm`, which is the seam and is supposed to call them, has two. --- .github/workflows/ci.yml | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 931fe95..985671f 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -274,8 +274,14 @@ jobs: # examples/09-heterogeneous has. grep -q '^import island_interface.kernels;' src/app.cppm \ || { echo "FAIL: the seam no longer imports the generated module"; exit 1; } - if grep -q 'saxpy_device\|scale_device' src/main.cpp; then + # CODE, NOT TEXT. `src/main.cpp` explains in a comment that it names + # no boundary symbol, and a plain grep matches that sentence -- the + # criterion then reports the file for saying what it does. Comment + # lines are stripped first, which is the difference between asking + # "does this file mention the name" and "does this file use it". + if sed 's://.*::' src/main.cpp | grep -q 'saxpy_device\|scale_device'; then echo "FAIL: the consumer names a boundary symbol; the seam exists to hide it" + sed 's://.*::' src/main.cpp | grep -n 'saxpy_device\|scale_device' exit 1 fi # AND THE CHECK THAT INCLUDE USED TO DO IS STILL THERE. The compiler From a6f5c1821980c8219473a464ab5444d9abafa3c7 Mon Sep 17 00:00:00 2001 From: sunrisepeak Date: Tue, 8 Sep 2026 02:24:33 +0800 Subject: [PATCH 20/20] ci: two more criteria that read text where they meant code `spirv-zero-config` explains in its own header that `host-module = true` is not written and not needed, and the step asserting that greps the manifest -- so the fixture was reported for saying what it does. Comment lines are stripped first. The island step's `grep 'include' build.mcpp` had the mirror-image problem: it matched that program's comments and its `mcpp::include_dir` line, so it would have passed with the forced include removed. A criterion that cannot fail is not one. It names `force_include_flags` now, which is the mechanism it is there to protect. Both were found by reading every CI assertion that greps a source file rather than by waiting for the next round trip; the other sixteen are either over program output, which has no comments, or over patterns no comment in those files carries. --- .github/workflows/ci.yml | 16 ++++++++++++++-- 1 file changed, 14 insertions(+), 2 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 985671f..254c727 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -288,7 +288,11 @@ jobs: # sees the declarations, so a definition whose signature drifted from # the generated header fails where it was written rather than at the # link. - grep -q 'include' build.mcpp \ + # THE MECHANISM, NOT THE WORD. `grep 'include'` matched this build + # program's own comments and its `mcpp::include_dir` line, so it + # would have passed with the forced include removed -- a criterion + # that cannot fail is not one. + grep -v '^[[:space:]]*//' build.mcpp | grep -q 'force_include_flags' \ || { echo "FAIL: the build program no longer forces the header in"; exit 1; } # THE MARKER SELECTS, AND THE UNMARKED MUST NOT TRAVEL. The island has # a static helper and an unmarked external function; a scan that @@ -392,7 +396,15 @@ jobs: working-directory: tests/spirv-zero-config run: | test ! -f build.mcpp || { echo "FAIL: the fixture has its own build program"; exit 1; } - grep -q 'host-module' mcpp.toml && { echo "FAIL: the fixture declares host-module"; exit 1; } + # STRIPPED OF COMMENTS, for the same reason the island step is: this + # fixture EXPLAINS in its own header that `host-module = true` is not + # written and not needed, so a plain grep reports the file for saying + # what it does. The property is about the manifest's keys. + if grep -v '^[[:space:]]*#' mcpp.toml | grep -q 'host-module'; then + echo "FAIL: the fixture declares host-module" + grep -v '^[[:space:]]*#' mcpp.toml | grep -n 'host-module' + exit 1 + fi "$MCPP" build | tee build.log grep -q 'Rules mcpp.rules.spirv' build.log \ || { echo "FAIL: the build did not report which rule it ran"; exit 1; }