diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 0633fd8..254c727 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` @@ -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,68 @@ 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 '^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" + exit 1 + fi + echo "ok: the consumer names only the module" + # And what it imports really is a generated module interface. + # + # `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 <' "$iface"; 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 @@ -98,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 @@ -116,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 || { @@ -139,6 +227,278 @@ 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.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 + # 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 + # 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/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; } + # 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 + # sees the declarations, so a definition whose signature drifted from + # the generated header fails where it was written rather than at the + # link. + # 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 + # 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" + + # 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`, + # 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; } + # 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; } + "$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 + # 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 +514,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. @@ -208,11 +594,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)" @@ -492,11 +891,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)" diff --git a/README.md b/README.md index 06c56af..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,9 +49,11 @@ 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` | 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 | +| `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 @@ -59,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 @@ -112,6 +114,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 | @@ -124,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 @@ -134,10 +144,287 @@ 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. +## 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 }; +``` + +### 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 +`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. + +**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 +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 +now, and every consumer reaches the same array through the accessor. + +## `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.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: + +```c +// 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) { ... } +``` + +```cpp +// build.mcpp +mcpp::tools::island::options opt; +opt.module_name = "myapp.kernels"; +opt.out_dir = std::string(mcpp::out_dir()) + "/island"; + +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()); +``` + +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 +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 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: + +``` +src/kernels/saxpy.c:22:5: error: conflicting types for 'scale_device' +``` + +**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 +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 @@ -151,7 +438,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..85c431e 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"] @@ -19,14 +19,76 @@ 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-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 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+) ────────────────────── # @@ -139,14 +201,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) @@ -157,16 +220,33 @@ 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. +# +# `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. +# +# 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" # 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/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/slang.cppm b/rules/slang.cppm new file mode 100644 index 0000000..932be07 --- /dev/null +++ b/rules/slang.cppm @@ -0,0 +1,467 @@ +// 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; + // 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; +} + +// 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"; + // 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; + + 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..4880587 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,42 @@ 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(); + + // 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 + // 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 +418,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 +430,61 @@ 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; + // 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; +} + // 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 +635,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 +664,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 +710,57 @@ 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"; + // 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"; + // 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; @@ -606,10 +773,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(); @@ -637,16 +809,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()); @@ -654,7 +832,40 @@ 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.store = opt.storage; + 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()); + // 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 + // 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/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 diff --git a/src/plugins.cppm b/src/plugins.cppm index dbb427f..344df85 100644 --- a/src/plugins.cppm +++ b/src/plugins.cppm @@ -1,16 +1,732 @@ -// 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 }; + +// 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`. + 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. + // 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 `. + // + // 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_; + 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 + // 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 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 + // 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; +} + +// 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"; +} + +// `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); + // 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", + identifier(it.identifier, "payload"), base, base); + } + reopen({}); + 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 +// 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); + // 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; + + // 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; + 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 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); + 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(); + if (!write_if_different(out.impl_file, impl)) { + 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; +} + +// 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() { + // 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(); + return identifier(leaf, "app"); +} + +} // namespace mcpp::plugins::surface diff --git a/tests/all-rules-compile/build.mcpp b/tests/all-rules-compile/build.mcpp index f265633..120dbd1 100644 --- a/tests/all-rules-compile/build.mcpp +++ b/tests/all-rules-compile/build.mcpp @@ -8,9 +8,11 @@ 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; +import mcpp.tools.island; int main() { // No accelerator is named, so each of these returns true without looking @@ -19,6 +21,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 @@ -26,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 4c3f3e4..78e8c6b 100644 --- a/tests/all-rules-compile/mcpp.toml +++ b/tests/all-rules-compile/mcpp.toml @@ -36,8 +36,9 @@ 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", - "tools-embed", + "rules-ascendc", "rules-cuda", "rules-hip", "rules-slang", "rules-spirv", + "rules-sycl", + "tools-embed", "tools-island", ], host-module = true } # NO `accel`, and that is the whole design: with none, every rule returns 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/island-interface/build.mcpp b/tests/island-interface/build.mcpp new file mode 100644 index 0000000..1466d49 --- /dev/null +++ b/tests/island-interface/build.mcpp @@ -0,0 +1,61 @@ +import std; +import mcpp; +import mcpp.tools.island; + +// THE ENTRY POINTS ARE READ FROM WHERE THEY ARE DEFINED, AND THE ISLAND WRITES +// NOTHING ELSE. +// +// `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 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() { + // 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()); + + 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::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()); + return 1; + } + + 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::tools::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/mcpp.toml b/tests/island-interface/mcpp.toml new file mode 100644 index 0000000..c41d526 --- /dev/null +++ b/tests/island-interface/mcpp.toml @@ -0,0 +1,52 @@ +# The consumer that tests `mcpp.tools.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" +accelerators = ["vulkan"] + +[language] +standard = "c++23" +modules = true +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] +accel = "vulkan1.2" +sources = ["src/*.cppm", "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" +main = "src/main.cpp" 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/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/tests/island-interface/src/kernels/saxpy.c b/tests/island-interface/src/kernels/saxpy.c new file mode 100644 index 0000000..fe55b0a --- /dev/null +++ b/tests/island-interface/src/kernels/saxpy.c @@ -0,0 +1,33 @@ +/* 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, 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. */ + +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; +} + +/* 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); } diff --git a/tests/island-interface/src/main.cpp b/tests/island-interface/src/main.cpp new file mode 100644 index 0000000..ab85a28 --- /dev/null +++ b/tests/island-interface/src/main.cpp @@ -0,0 +1,38 @@ +// 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.app; + +int main() { + const std::vector x{1, 2, 3, 4}; + const std::vector y{10, 20, 30, 40}; + + auto out = island_interface::saxpy(2.0f, x, y); + if (!out) { + std::cout << "BAD: saxpy failed\n"; + return 1; + } + 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 = 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 + // only the first would still satisfy a single-function fixture. + 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..1101b8c --- /dev/null +++ b/tests/spirv-module-consumer/mcpp.toml @@ -0,0 +1,51 @@ +# 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 +# `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. +# +# 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 = "shader-app" +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.shader-app] +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/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 new file mode 100644 index 0000000..995d418 --- /dev/null +++ b/tests/spirv-module-consumer/src/main.cpp @@ -0,0 +1,56 @@ +// 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 shader_app.shaders; + +namespace shaders = shader_app::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()); + // 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 + // 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/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; +} 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; +} diff --git a/tools/embed.cppm b/tools/embed.cppm index 6a1dd82..4e7a85a 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. @@ -84,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() { @@ -215,4 +218,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 diff --git a/tools/island.cppm b/tools/island.cppm new file mode 100644 index 0000000..5af24ca --- /dev/null +++ b/tools/island.cppm @@ -0,0 +1,359 @@ +// 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. +// +// 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) { + 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()) 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; +} + +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