From 3f6f75fc172074157ab8ab28171651939c8538be Mon Sep 17 00:00:00 2001 From: speak-agent <248744407+speak-agent@users.noreply.github.com> Date: Tue, 8 Sep 2026 01:33:37 +0800 Subject: [PATCH] examples: the module surface, and one header kept for contrast The payload lane and the boundary lane are separate, and each example now shows one of them. `vulkan` reaches its SPIR-V through a MODULE. `src/vulkan/saxpy.cpp` opened with `#include "scale_comp.h"` -- a name no line in the project produced and no reader could derive without opening the rule -- and now imports `vulkan_saxpy.shaders`, whose name and whose `scale_comp()` accessor are both derived from names the project already wrote. The accessor answers with the address and the byte count together, which is what makes `sizeof` the wrong question rather than merely an awkward one: the rule may put the words in an object, and there is then no array to take the size of. `cuda` and `sycl` GENERATE their `extern "C"` boundary. Each entry point is marked with `MCPP_EXPORT_C` where it is defined, and `mcpp.tools.island` writes the header the island's compiler reads and the module the seam imports. `include/saxpy/saxpy.h` is gone from both. Both halves are handed to the scan, so a signature that drifted between the island and the CPU fallback is refused while the boundary is generated -- the one point at which both texts exist at once, and the only check available at a boundary where C language linkage does not mangle and the two halves are never in one link. The island names nothing and the host half writes one `#include`, and that asymmetry is measured rather than incidental. The island is compiled by a driver mcpp did not invoke, so its rule takes forced-include flags for that one command line. The host half is ordinary project C++, and the only project-wide channel would force the header into every C++ translation unit including the seam, where declarations ahead of `export module` are ill-formed. `hip` keeps its hand-written header, and its README says why. It computes the same thing through the same seam as `cuda`, so the difference between the two examples is exactly this one thing, and the hand-written form is what every C boundary looks like today. Neither is deprecated; the generated form is the default because the copy it removes is the one that goes wrong quietly. `multi-backend` reaches its SPIR-V as `opkit.shaders` too, and keeps `include/opkit/opkit.h` written by hand -- which is the distinction this example is for. That header is a LIBRARY boundary: consumers outside this tree compile against it, so it is not an intermediate and generating it would not remove a copy, it would move one. The shader payload is an intermediate, and it moves. `examples/10-graphics/offscreen` moves too, and it is the one example CI both builds AND RUNS on two platforms -- so the surface is exercised there rather than merely compiled. Two includes become one import, and the two stages are `offscreen::shaders::triangle_vert()` and `::triangle_frag()`. `cann` is unchanged. Its boundary has a different shape -- `saxpy_device` between the seam and a host glue file, `saxpy_launch` between that glue and the kernel -- and both halves of the first one are compiled by mcpp, so the forced-include asymmetry above does not apply. Its README already records that the example does not build yet. Verified on this host: cuda and sycl build and run on an RTX 4080 with the generated boundary on the island's own command line, vulkan runs on the same device with its shader reached as a module, multi-backend runs on lavapipe with `--accel "vulkan1.2"`, the offscreen triangle renders on lavapipe at centre pixel `(124, 70, 62, 255)` -- the same pixel its CPU leg produces, which is the assertion CI makes -- and each CPU leg builds and runs under `--no-accel`. `docs/20`'s two-lanes section said the seam is written by hand and the shader lane's equivalent is generated. Half of that still holds: the seam is a design decision. What sits under it is each signature written a second time, and that half is now generated too. The section says so, and gains the table that maps a file name to the call it arrives as -- which is the question a reader of that document actually has. Every example under `examples/09-heterogeneous` and `examples/10-graphics` pins 0.3.0 now, `hip` and `cann` included. Neither uses anything new, and leaving them behind would be defensible in isolation -- but `hip` exists to be read beside `cuda`, and a reader comparing the two would find a second difference the text never explains. The floor sentences move with them. The six plugin pins in `docs/05` and `docs/20` move too, in both languages. A version in a document is the line someone pastes, and 0.2.1 and 0.2.4 predate every declaration those sections now describe -- a reader who copied one would get a build in which the rules never route. The Windows graphics run step was brought back as a MEASUREMENT and is withdrawn again, this time with an answer. Four rounds established what it is not -- a missing `vulkan-1.dll`, an unparseable ICD manifest (a real defect, fixed in openxlings/xim-pkgindex#781, and not this failure's cause), a failed `LoadLibrary`, or the driver-files variable's vintage -- and the last round asked the loader instead of guessing: INFO: Loader is running with elevated permissions. Environment variable VK_DRIVER_FILES will be ignored INFO: Loader is running with elevated permissions. Environment variable VK_ICD_FILENAMES will be ignored ERROR | DRIVER: Registry lookup failed to get ICD manifest files. A GitHub Windows runner runs elevated, and the loader discards every driver-path environment variable when it is -- it will not let a path a non-administrator could write inject a driver into an elevated process. It falls back to the registry, which has no ICD. No package change reaches that; the mechanism this example uses on Linux and macOS is unavailable to an elevated process on Windows, and reaching lavapipe there means a registry entry under `HKEY_LOCAL_MACHINE\SOFTWARE\Khronos\Vulkan\Drivers` or an unelevated run. The four eliminations and the loader's own words are recorded beside the withdrawn step, so the next attempt starts from there rather than from the first hypothesis. Windows keeps building the Vulkan half and running the CPU fallback. Requires mcpp:plugins 0.3.0. --- ...9-07-module-first-heterogeneous-surface.md | 49 +++++++++++ .github/workflows/ci-windows.yml | 81 ++++++++++++------- docs/01-examples.md | 8 +- docs/05-mcpp-toml.md | 4 +- docs/20-heterogeneous-builds.md | 48 +++++++++-- docs/zh/01-examples.md | 8 +- docs/zh/05-mcpp-toml.md | 4 +- docs/zh/20-heterogeneous-builds.md | 38 +++++++-- examples/09-heterogeneous/README.md | 2 +- examples/09-heterogeneous/cann/app/mcpp.toml | 4 +- examples/09-heterogeneous/cuda/README.md | 21 ++++- examples/09-heterogeneous/cuda/app/build.mcpp | 70 +++++++++++++++- .../cuda/app/include/saxpy/saxpy.h | 31 ------- examples/09-heterogeneous/cuda/app/mcpp.toml | 10 ++- .../09-heterogeneous/cuda/app/src/app.cppm | 7 +- .../cuda/app/src/cpu/saxpy.cpp | 23 +++++- .../cuda/app/src/kernels/saxpy.cu | 9 ++- examples/09-heterogeneous/hip/README.md | 29 ++++++- examples/09-heterogeneous/hip/app/mcpp.toml | 4 +- .../09-heterogeneous/multi-backend/README.md | 2 +- .../09-heterogeneous/multi-backend/build.mcpp | 9 +++ .../09-heterogeneous/multi-backend/mcpp.toml | 2 +- .../src/backends/vulkan/host.cpp | 17 +++- examples/09-heterogeneous/sycl/README.md | 5 +- examples/09-heterogeneous/sycl/app/build.mcpp | 54 ++++++++++++- .../sycl/app/include/saxpy/saxpy.h | 31 ------- examples/09-heterogeneous/sycl/app/mcpp.toml | 8 +- .../09-heterogeneous/sycl/app/src/app.cppm | 7 +- .../sycl/app/src/cpu/saxpy.cpp | 23 +++++- .../sycl/app/src/kernels/saxpy.sycl | 9 ++- examples/09-heterogeneous/vulkan/README.md | 5 +- .../09-heterogeneous/vulkan/app/build.mcpp | 15 ++++ .../09-heterogeneous/vulkan/app/mcpp.toml | 2 +- .../vulkan/app/src/vulkan/saxpy.cpp | 36 ++++++--- examples/10-graphics/README.md | 17 ++-- examples/10-graphics/offscreen/build.mcpp | 16 +++- examples/10-graphics/offscreen/mcpp.toml | 21 +++-- .../offscreen/src/vulkan/render.cpp | 18 ++++- 38 files changed, 548 insertions(+), 199 deletions(-) delete mode 100644 examples/09-heterogeneous/cuda/app/include/saxpy/saxpy.h delete mode 100644 examples/09-heterogeneous/sycl/app/include/saxpy/saxpy.h diff --git a/.agents/docs/2026-09-07-module-first-heterogeneous-surface.md b/.agents/docs/2026-09-07-module-first-heterogeneous-surface.md index 46aca9e20..d48b1f76c 100644 --- a/.agents/docs/2026-09-07-module-first-heterogeneous-surface.md +++ b/.agents/docs/2026-09-07-module-first-heterogeneous-surface.md @@ -846,6 +846,55 @@ and a newer one moves a project to the module surface with nothing declared. | Seamless upgrade | A project that writes nothing new keeps its behaviour | Header storage stays the default under `modules = false`; the default surface follows `[language] modules` | | Test coverage | Each surface and each storage has a consumer that runs | P7 | +### 13.3.1 What the self-review found after the surface was working + +Four findings, three of them in names the surface generates. Every one was found +by giving the generator an input nobody had tried, and none of them by reading +the code -- which is the reason they are recorded together. + +**The user-facing name came from the wrong question.** The module a consumer +imports was derived from the leaf of `MCPP_MANIFEST_DIR` -- a directory name -- +because nothing in the build-program contract answered "what package am I +building". `examples/09-heterogeneous/vulkan/app/` declares +`name = "vulkan-saxpy"` and generated `app.shaders`. Worse than the defect is +that no fixture could see it: all fourteen had a package name equal to their +directory name, so both derivations produced the same string. Fixed by +`MCPP_PKG_NAME` in the engine, and by making one fixture's two names differ on +purpose. + +**A generated name may not be a C++ keyword.** Six sites ran the same character +filter and none checked for reserved words. `shaders/default/` produced +`namespace default {` in a generated file. Fixed by one `surface::identifier` +that owns all three transformations, applied at the point where the accessor is +EMITTED rather than in each producer -- a rule that builds the name from a file +stem cannot know it has produced a keyword until it reaches the line that writes +the function's name. + +**A seam's two halves are compared nowhere else.** The island and the host +fallback define one `extern "C"` boundary and are never in one link, and C +language linkage does not mangle, so two that declare a name differently build +cleanly and the artifact reads its arguments by whichever it was compiled with. +`mcpp.tools.island::scan` is handed both files and is therefore the only point +at which both texts exist at once. It refuses there. + +**One fix was written and withdrawn.** `host-module` is inferred from a rule +feature the consumer REQUESTS. A rule package whose rule sits in its own +`[features] default` activates without being named, so the rule modules are +collected -- that reads the resolved set -- while the inference reads the +requested one. Extending the inference to the dependency's defaults was measured +against a probe package and made the outcome WORSE: the refusal went from + + error: build.mcpp imports 'probe.rules.probe', and no dependency provides + it as a host module. + declared without `host-module = true`: probe.rules + +to the same refusal plus `importable here: rules`, which says the dependency is +wired up while the feature's own module still is not. The gap is further down, +in when a default feature's sources are folded into the set `units()` +enumerates. `mcpp:plugins` declares `default = []`, so nothing shipped reaches +it, and the existing refusal already names exactly what to add. Recorded here +rather than half-fixed. + ### 13.4 Release and verification sequence 1. mcpp pull request: E1 through E5. CI green on the head, then merged, then the diff --git a/.github/workflows/ci-windows.yml b/.github/workflows/ci-windows.yml index dc11409e5..f9d0d38d9 100644 --- a/.github/workflows/ci-windows.yml +++ b/.github/workflows/ci-windows.yml @@ -456,41 +456,62 @@ jobs: done echo "ok: both shader stages compiled and the Vulkan half linked" - # WINDOWS STAYS AT "BUILDS", AND WHAT REMAINS OPEN IS THE DRIVER. + # WINDOWS STAYS AT "BUILDS", AND FOUR MEASUREMENTS SAY WHY IT IS NOT + # WAITING ON ANY OF THE THINGS IT WAS THOUGHT TO BE. # - # A run step was written, pushed, and measured (run 34135108981). The - # manifest declaration worked -- `Provisioning [xlings.workspace] entries - # (xim:mesa-lavapipe@26.2.0)` -- the ICD manifest was found in the store, - # and the program printed `render unavailable`, which is `src/main.cpp` - # reporting that the render function returned nothing. + # 1. Not a missing `vulkan-1.dll`. The program printed `render + # unavailable`, which `src/main.cpp` writes after the render function + # returns nothing. A process that could not resolve `vkCreateInstance` + # from that DLL fails during image load and prints nothing at all. + # mcpp-index's own `vulkan-tests` member calls + # `vkEnumerateInstanceVersion` on the windows shards and passes. # - # THAT OUTPUT PLACES THE FAILURE AFTER THE LOADER, NOT AT IT. The Vulkan - # leg imports `vkCreateInstance` from `vulkan-1.dll` through the import - # library, so a process missing that DLL fails during image load and - # prints nothing at all. It printed. The loader was there, it ran, and it - # enumerated no device -- which is a statement about the ICD. + # 2. Not an unparseable ICD manifest, though that WAS a real defect. + # `xim:mesa-lavapipe` wrote its rewritten `library_path` into the + # JSON string unescaped, so `C:\Users\...` carried `\U` and the + # loader's cJSON parser rejected the file -- an ICD it skips with no + # error. Fixed in openxlings/xim-pkgindex#781. With the fix the runner + # reads + # "library_path": "C:/Users/.../lib/vulkan_lvp.dll" + # and the manifest parses. The program still prints `render + # unavailable`. # - # An earlier revision of this comment named a missing `vulkan-1.dll` as - # the cause. That was an inference from the descriptor's note about static - # linkage rather than a reading of the failure, and the log refutes it. - # Two facts stand against it: this program ran, and mcpp-index's own - # `vulkan-tests` member calls `vkEnumerateInstanceVersion` on the windows - # shards and passes. + # 3. Not a failed `LoadLibrary`. `vulkan_lvp.dll` imports only + # ADVAPI32, GDI32, KERNEL32, ntdll, ole32, SHELL32 and USER32, all of + # which any Windows has; the payload ships nothing else that could be + # missing. # - # A separate measurement, recorded here because it removes a second thing - # from the list of suspects: the Khronos loader in `compat:vulkan` builds - # into a working `vulkan-1.dll` from the source the index already carries - # (mingw, 265 exports matching upstream's `vulkan-1.def` exactly, DllMain - # present, importing only ADVAPI32/CFGMGR32/KERNEL32/msvcrt). The - # descriptor's note argues that a Windows loader must be a DLL, not that - # it cannot be built -- so a hermetic Windows loader is available whenever - # it is wanted. It is not what this step is waiting on. + # 4. Not the environment variable's vintage. Both `VK_DRIVER_FILES` and + # `VK_ICD_FILENAMES` were set, which covers loaders on either side of + # 1.3.234. # - # What this step is waiting on is `xim:mesa-lavapipe`'s Windows payload - # producing a device under an mcpp-launched process. Until that is - # measured, the job builds the Vulkan half and runs the CPU fallback. - # macOS is not in the same position: MoltenVK enumerates once the instance - # asks for portability, which `src/vulkan/render.cpp` now does. + # What the loader itself said under `VK_LOADER_DEBUG=all`: + # + # INFO: Loader is running with elevated permissions. + # Environment variable VK_DRIVER_FILES will be ignored + # INFO: Loader is running with elevated permissions. + # Environment variable VK_ICD_FILENAMES will be ignored + # DRIVER: Found no registry files in + # HKEY_LOCAL_MACHINE\SOFTWARE\Khronos\Vulkan\Drivers + # ERROR | DRIVER: windows_read_data_files_in_registry: Registry lookup + # failed to get ICD manifest files. Possibly missing Vulkan driver? + # + # THAT IS THE WHOLE ANSWER, AND NO PACKAGE CHANGE REACHES IT. A GitHub + # Windows runner runs elevated, and the loader discards every driver-path + # environment variable when it is -- it will not let a path a + # non-administrator could write inject a driver into an elevated process. + # It then falls back to the registry, which has no ICD. The mechanism this + # example uses on Linux and macOS is simply unavailable here. + # + # Reaching lavapipe on Windows therefore means registering the ICD under + # `HKEY_LOCAL_MACHINE\SOFTWARE\Khronos\Vulkan\Drivers`, or running the + # program unelevated. `xim:mesa-lavapipe`'s `config()` says outright that + # it places the payload and leaves naming the ICD to the consumer, so the + # registry entry is a decision for that package or for this repository, + # not a defect in either + # + # So the platform builds the Vulkan half and runs the CPU fallback, and + # the next attempt starts from here rather than from the top. - name: "Toolchain: LLVM — build mcpp (self-host)" shell: bash diff --git a/docs/01-examples.md b/docs/01-examples.md index 98285aec6..33edb693b 100644 --- a/docs/01-examples.md +++ b/docs/01-examples.md @@ -32,10 +32,10 @@ examples. | 07 | [`examples/07-project-subos`](../examples/07-project-subos/) | A build program that finds its tools in the environment the project declared | `[xlings] subos`, `[xlings.workspace]`, a build program whose `PATH` is the environment the project named | | 08 | [`examples/08-build-rules`](../examples/08-build-rules/) | Two rule packages and a project that uses both | `host-module = true`, `[build-dependencies]`, `mcpp::action` with `role = "check"` | | 09 | [`examples/09-heterogeneous`](../examples/09-heterogeneous/) | One computation on a device, in several programming models, with a CPU fallback in each; plus one artifact carrying several backends at once | `accel`, constrained source globs, the seam module, rule packages from `mcpp:plugins`, `cfg(accelerator = …)` | -| 09a | [`…/cuda`](../examples/09-heterogeneous/cuda/) | A CUDA kernel behind a seam module | `mcpp.rules.cuda`, `mcpp::action` with `role = "object"`, the driver stated as a fact and a floor | -| 09b | [`…/vulkan`](../examples/09-heterogeneous/vulkan/) | The same computation as a Vulkan compute shader, on a GPU or on the CPU | `mcpp.rules.spirv`, `mcpp::action` with `role = "source"`, generated headers, a software driver as a payload | -| 09c | [`…/sycl`](../examples/09-heterogeneous/sycl/) | The same computation as a SYCL kernel, compiled by a second compiler | `mcpp.rules.sycl`, the `.sycl` device extension, a chained `mcpp::action` for the device link, `compat:sycl-runtime` | -| 09d | [`…/hip`](../examples/09-heterogeneous/hip/) | The same computation in HIP, reaching an NVIDIA device | `mcpp.rules.hip`, HIP as a header layer over the CUDA runtime, a two-chunk `accel` | +| 09a | [`…/cuda`](../examples/09-heterogeneous/cuda/) | A CUDA kernel behind a seam module, with the `extern "C"` boundary GENERATED | `mcpp.rules.cuda`, `mcpp.tools.island`, `mcpp::action` with `role = "object"`, the driver stated as a fact and a floor | +| 09b | [`…/vulkan`](../examples/09-heterogeneous/vulkan/) | The same computation as a Vulkan compute shader, with the SPIR-V payload reached as a MODULE | `mcpp.rules.spirv`, the module surface, `mcpp::action` with `role = "source"`, a software driver as a payload | +| 09c | [`…/sycl`](../examples/09-heterogeneous/sycl/) | The same computation as a SYCL kernel, compiled by a second compiler | `mcpp.rules.sycl`, `mcpp.tools.island`, the `.sycl` device extension, a chained `mcpp::action` for the device link, `compat:sycl-runtime` | +| 09d | [`…/hip`](../examples/09-heterogeneous/hip/) | The same computation in HIP, with the boundary WRITTEN BY HAND — the contrast against 09a | `mcpp.rules.hip`, HIP as a header layer over the CUDA runtime, a two-chunk `accel` | | 09e | [`…/multi-backend`](../examples/09-heterogeneous/multi-backend/) | Several backends in ONE artifact, chosen at run time — the library shape, not the program shape | `accel` as a set, `cfg(accelerator = "none")` and its negation, a dispatch chain, a module seam over a C island boundary | | 09f | [`…/cann`](../examples/09-heterogeneous/cann/) | An Ascend C kernel behind the same seam. **Does not build yet** — its README names the two missing pieces | the `.asc` device extension, `op_kernel`/`op_host` as an island CANN already has, `accelerator = "none"` for the fallback | | 10 | [`examples/10-graphics`](../examples/10-graphics/) | Graphics rather than compute: a rendering pipeline whose result is pixels | `mcpp.rules.spirv` for the vertex and fragment stages, offscreen rendering as the assertable form | diff --git a/docs/05-mcpp-toml.md b/docs/05-mcpp-toml.md index bbe6c2c5b..f789ff7ab 100644 --- a/docs/05-mcpp-toml.md +++ b/docs/05-mcpp-toml.md @@ -2308,7 +2308,7 @@ have none, and nothing is what they write: ```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 } ``` That one edge is the whole declaration. The rule package names the packages its @@ -2546,7 +2546,7 @@ rules-spirv = { sources = ["rules/spirv.cppm"] } # export module mcpp.rules.spir ```toml # a consumer [build-dependencies.mcpp] -plugins = { version = "0.2.1", features = ["rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-spirv"], host-module = true } ``` **`[build-dependencies]`, not `[dependencies]`** — a rule package is the case diff --git a/docs/20-heterogeneous-builds.md b/docs/20-heterogeneous-builds.md index d47299a15..204190876 100644 --- a/docs/20-heterogeneous-builds.md +++ b/docs/20-heterogeneous-builds.md @@ -185,13 +185,20 @@ could read a BMI. The seam above is written by hand, and the shader lane's equivalent is generated. That is not an inconsistency; the two lanes carry different things. -**A device translation unit is code.** Its interface is a design decision -- -which functions, which types, what happens on failure -- and no generator can -make that decision well. So CUDA, HIP, SYCL and Ascend C get a hand-written -seam, and the `extern "C"` header exists for the ABI reason above. Both are -already invisible to a consumer: only the seam includes the header, and -everything downstream writes `import app.saxpy`. **These lanes are module-first -today and always have been.** +**A device translation unit is code, and the seam over it is a design +decision** -- which functions, which types, what happens on failure. No +generator makes that decision well, so CUDA, HIP, SYCL and Ascend C get a +hand-written seam, and everything downstream writes `import app.saxpy`. + +**The `extern "C"` boundary UNDER that seam is not a design decision.** It is +each entry point's signature, stated a second time in a header, at the one place +where a disagreement is invisible: C language linkage does not mangle, and a +device island and its host fallback are never in one link. `mcpp.tools.island` +reads the marked declarations out of both implementations and writes that header +and a module over it, so the signatures exist once and two halves that disagree +are refused where both texts are in front of the generator. +`examples/09-heterogeneous/cuda` and `.../sycl` are that shape; +`.../hip` keeps its header written by hand, so the two can be read side by side. **A shader or an embedded file is data.** Its interface is an address and a size, which is mechanical, so a rule package generates it and a consumer writes @@ -203,6 +210,31 @@ So the rule is not "generate the interface" or "write it by hand". It is: a mechanical interface is generated, a designed one is written, and in both cases the header is an intermediate that no consumer names. +#### The name a payload arrives under + +The module and the namespace are one identifier path, derived from names the +project already wrote. + +| Written | Reached as | +|---|---| +| `[package] name = "myapp"` | module root `myapp` | +| `shaders/scale.comp` | `myapp::shaders::scale_comp()` | +| `shaders/a/scale.comp` | `myapp::shaders::a::scale_comp()` | +| a `MCPP_EXPORT_C` entry point | `export using ::the_name;` in the boundary module | + +The root is the **package's** name with non-identifier characters replaced, not +its directory's -- the two differ whenever a package sits under a generic folder, +and mcpp reports the package name to a build program from 2026.9.7.1 for exactly +this. The stem and the stage make the accessor (`scale.comp` -> `scale_comp`), +and the directory below the group's base becomes namespace segments, which is +what makes two shaders sharing a stem two things rather than a collision. A +project that wants another name passes one; `examples/09-heterogeneous/cuda` +does, so its boundary is `app.kernels` beside its seam `app.saxpy`. + +The accessor answers with the address and the byte count together. `sizeof` is +not merely awkward at this boundary, it is unanswerable: the words may be in an +object rather than in an array, and there is then nothing to take the size of. + ## Compiling an island The command that invokes a device compiler is not built into mcpp. It is @@ -221,7 +253,7 @@ A project that wants a CUDA island writes one edge: ```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 } ``` and nothing else. The vendor toolkit, its runtime and whatever else the rule diff --git a/docs/zh/01-examples.md b/docs/zh/01-examples.md index 46ac445c7..ae8179747 100644 --- a/docs/zh/01-examples.md +++ b/docs/zh/01-examples.md @@ -29,10 +29,10 @@ mcpp build && mcpp run | 07 | [`examples/07-project-subos`](../../examples/07-project-subos/) | 构建程序在工程声明的环境里找工具,而不是问机器上恰好有什么 | `[xlings] subos`、`[xlings.workspace]`、构建程序的 `PATH` 来自工程声明的那个环境 | | 08 | [`examples/08-build-rules`](../../examples/08-build-rules/) | 两个规则包,以及同时用到它们的工程 | `host-module = true`、`[build-dependencies]`、`role = "check"` 的 `mcpp::action` | | 09 | [`examples/09-heterogeneous`](../../examples/09-heterogeneous/) | 同一个计算在设备上跑,写成多种编程模型,每种都带 CPU 回退;外加一个同时携带多个后端的产物 | `accel`、带约束的 source glob、接缝模块、来自 `mcpp:plugins` 的规则包、`cfg(accelerator = …)` | -| 09a | [`…/cuda`](../../examples/09-heterogeneous/cuda/) | 接缝模块背后的 CUDA kernel | `mcpp.rules.cuda`、`role = "object"` 的 `mcpp::action`、把驱动陈述为 fact 与 floor | -| 09b | [`…/vulkan`](../../examples/09-heterogeneous/vulkan/) | 同一个计算写成 Vulkan compute shader,在 GPU 上或在 CPU 上 | `mcpp.rules.spirv`、`role = "source"` 的 `mcpp::action`、生成的头文件、作为载荷的软件驱动 | -| 09c | [`…/sycl`](../../examples/09-heterogeneous/sycl/) | 同一个计算写成 SYCL kernel,由第二个编译器编译 | `mcpp.rules.sycl`、`.sycl` 设备扩展名、为 device link 串起来的 `mcpp::action`、`compat:sycl-runtime` | -| 09d | [`…/hip`](../../examples/09-heterogeneous/hip/) | 同一个计算写成 HIP,够到一台 NVIDIA 设备 | `mcpp.rules.hip`、HIP 作为 CUDA 运行时之上的一层头文件、两段式的 `accel` | +| 09a | [`…/cuda`](../../examples/09-heterogeneous/cuda/) | 接缝模块背后的 CUDA kernel,`extern "C"` 边界是**生成**的 | `mcpp.rules.cuda`、`mcpp.tools.island`、`role = "object"` 的 `mcpp::action`、把驱动陈述为 fact 与 floor | +| 09b | [`…/vulkan`](../../examples/09-heterogeneous/vulkan/) | 同一个计算写成 Vulkan compute shader,SPIR-V 载荷以**模块**到达 | `mcpp.rules.spirv`、模块表面、`role = "source"` 的 `mcpp::action`、作为载荷的软件驱动 | +| 09c | [`…/sycl`](../../examples/09-heterogeneous/sycl/) | 同一个计算写成 SYCL kernel,由第二个编译器编译 | `mcpp.rules.sycl`、`mcpp.tools.island`、`.sycl` 设备扩展名、为 device link 串起来的 `mcpp::action`、`compat:sycl-runtime` | +| 09d | [`…/hip`](../../examples/09-heterogeneous/hip/) | 同一个计算写成 HIP,边界是**手写**的 —— 与 09a 的对照 | `mcpp.rules.hip`、HIP 作为 CUDA 运行时之上的一层头文件、两段式的 `accel` | | 09e | [`…/multi-backend`](../../examples/09-heterogeneous/multi-backend/) | 多个后端进**同一个产物**,运行期选择 —— 这是库的形态,不是程序的形态 | `accel` 作为集合、`cfg(accelerator = "none")` 及其否定、分发链、C 岛边界之上的模块接缝 | | 09f | [`…/cann`](../../examples/09-heterogeneous/cann/) | 同一道接缝背后的 Ascend C kernel。**目前还构建不了** —— README 里点明了缺的两块 | `.asc` 设备扩展名、CANN 本来就有的 `op_kernel`/`op_host` 岛、回退用 `accelerator = "none"` | | 10 | [`examples/10-graphics`](../../examples/10-graphics/) | 图形而不是计算:一条渲染管线,它的结果是像素 | `mcpp.rules.spirv` 编译顶点与片段两个阶段、离屏渲染作为可断言的形态 | diff --git a/docs/zh/05-mcpp-toml.md b/docs/zh/05-mcpp-toml.md index c360d2483..bdff7b08c 100644 --- a/docs/zh/05-mcpp-toml.md +++ b/docs/zh/05-mcpp-toml.md @@ -1964,7 +1964,7 @@ hardware = {} ```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 } ``` 这一条边就是全部声明。规则包在选中它的那个 feature、它所服务的加速器之下,声明自己 @@ -2166,7 +2166,7 @@ rules-spirv = { sources = ["rules/spirv.cppm"] } # export module mcpp.rules.spir ```toml # 消费者 [build-dependencies.mcpp] -plugins = { version = "0.2.1", features = ["rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-spirv"], host-module = true } ``` **用 `[build-dependencies]` 而不是 `[dependencies]`** —— 规则包正是 §2.6.1 描述的那种 diff --git a/docs/zh/20-heterogeneous-builds.md b/docs/zh/20-heterogeneous-builds.md index c889e77a0..f5654d77e 100644 --- a/docs/zh/20-heterogeneous-builds.md +++ b/docs/zh/20-heterogeneous-builds.md @@ -152,11 +152,16 @@ C 语言链接不做名字修饰,所以这两个是同一个符号。链接是 上面那个接缝是手写的,而 shader 那条 lane 的对应物是生成的。这不是不一致 —— 两条 lane 承载的东西不同。 -**设备编译单元是代码。** 它的接口是一个设计决定 —— 有哪些函数、什么类型、失败怎么报 -—— 而没有任何生成器能把这个决定做好。所以 CUDA、HIP、SYCL 与 Ascend C 有一个手写的 -接缝,而那个 `extern "C"` 头文件因为上面那条 ABI 理由而存在。两者对消费者**都已经是 -不可见的**:只有接缝包含那个头文件,下游一律写 `import app.saxpy`。**这些 lane 今天 -就是模块优先的,而且一直如此。** +**设备编译单元是代码,而它上面那层接缝是一个设计决定** —— 有哪些函数、什么类型、 +失败怎么报。没有任何生成器能把这个决定做好,所以 CUDA、HIP、SYCL 与 Ascend C 有一个 +手写的接缝,下游一律写 `import app.saxpy`。 + +**而接缝**底下**那个 `extern "C"` 边界不是设计决定。** 它是每个入口点的签名被第二次 +写出来,写在一个「不一致看不见」的位置上:C 语言链接不做名字修饰,而设备岛与它的宿主 +回落**永不同链**。`mcpp.tools.island` 从两个实现里读出被标记的声明,写出那个头文件和 +它之上的模块,于是签名只存在一次,而两个半边如果声明不一致,会在两份文本同时摆在生成器 +面前的那一刻被拒绝。`examples/09-heterogeneous/cuda` 与 `.../sycl` 是这个形状; +`.../hip` 保留手写的头文件,两者可以并排读。 **shader 或一份被嵌入的文件是数据。** 它的接口是一个地址加一个尺寸,这是机械的,所以 由规则包生成,消费者同样不写出任何生成物的名字,只写 `import myapp.shaders`。在 @@ -166,6 +171,27 @@ mcpp 2026.9.7.1 之前,那条 lane 是唯一的例外:生成的头文件**就是 所以规则不是「接口要生成」也不是「接口要手写」。规则是:**机械的接口生成,设计出来的 接口手写,而两种情况下头文件都是没有任何消费者会写出其名字的中间产物。** +#### 一份载荷以什么名字到达 + +模块名与命名空间是同一条标识符路径,由工程已经写下的名字推导。 + +| 写下的 | 到达时的样子 | +|---|---| +| `[package] name = "myapp"` | 模块根 `myapp` | +| `shaders/scale.comp` | `myapp::shaders::scale_comp()` | +| `shaders/a/scale.comp` | `myapp::shaders::a::scale_comp()` | +| 一个 `MCPP_EXPORT_C` 入口点 | 边界模块里的 `export using ::the_name;` | + +根取自**包名**(非标识符字符替换掉),不是目录名 —— 包放在一个通用目录下时两者不同, +而 mcpp 从 2026.9.7.1 起把包名报给构建程序,正是为了这一条。文件名的 stem 与 stage +构成访问器(`scale.comp` -> `scale_comp`),而相对该组基准目录的子目录成为命名空间段 —— +这就是让两个同 stem 的 shader 成为两样东西而不是一次冲突的机制。想要别的名字就传一个: +`examples/09-heterogeneous/cuda` 就传了,于是它的边界是 `app.kernels`,与接缝 +`app.saxpy` 同根。 + +访问器同时回答地址与字节数。在这个边界上 `sizeof` 不只是别扭,而是**答不出来**: +那些字可能放在一个对象里而不是一个数组里,那时根本没有数组可以取 size。 + ## 编译一个岛 调用设备编译器的那条命令不内置在 mcpp 里,而是由**构建规则包**提供 —— @@ -181,7 +207,7 @@ mcpp 2026.9.7.1 之前,那条 lane 是唯一的例外:生成的头文件**就是 ```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 } ``` **别的什么都不写。** 厂商工具包、它的运行时,以及规则需要的其余东西,都由**规则包** diff --git a/examples/09-heterogeneous/README.md b/examples/09-heterogeneous/README.md index 3d7ab36a4..88b9609f5 100644 --- a/examples/09-heterogeneous/README.md +++ b/examples/09-heterogeneous/README.md @@ -84,7 +84,7 @@ selects it and the accelerator it serves, so a project writes one edge and no ```toml [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-cuda"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-cuda", "tools-island"], host-module = true } ``` `multi-backend/` is the one example here that also pins a version, and it does diff --git a/examples/09-heterogeneous/cann/app/mcpp.toml b/examples/09-heterogeneous/cann/app/mcpp.toml index 01f57b5d0..37375fcc6 100644 --- a/examples/09-heterogeneous/cann/app/mcpp.toml +++ b/examples/09-heterogeneous/cann/app/mcpp.toml @@ -14,7 +14,7 @@ import_std = true # imports it as `mcpp.rules.ascendc`. `[build-dependencies]`, because a rule # package's library must never reach the target while its rule is wanted. [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-ascendc"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-ascendc"], host-module = true } # NO [xlings.workspace]. `mcpp.rules.ascendc` declares `xim:cann-toolkit` # itself, gated on this accelerator and on the feature that selects the rule -- @@ -29,7 +29,7 @@ plugins = { version = "0.2.4", features = ["rules-ascendc"], host-module = true # # The rule states `>=8.5.0` rather than an exact version: nothing here is # coupled to a driver, and that is the release the mixed-mode object and the -# simulators arrive in. Needs mcpp 2026.9.6.6 and mcpp:plugins 0.2.4. +# simulators arrive in. Needs mcpp 2026.9.7.1 and mcpp:plugins 0.3.0. [build] # `dav-2201` is the device architecture, the role `sm_89` plays for CUDA. The diff --git a/examples/09-heterogeneous/cuda/README.md b/examples/09-heterogeneous/cuda/README.md index 8c830dbdc..b8af616fd 100644 --- a/examples/09-heterogeneous/cuda/README.md +++ b/examples/09-heterogeneous/cuda/README.md @@ -10,7 +10,10 @@ app/ no BMI, and compiled only when the build asks for CUDA src/cpu/saxpy.cpp the same interface implemented for the host, compiled instead when it does not - include/saxpy/saxpy.h the island's interface: extern "C", no std types + (generated) the `extern "C"` boundary and the module over it, + written by `mcpp.tools.island` from the marked + declarations in the two implementations above -- + there is no header in this source tree src/app.cppm the seam: a module that turns the C interface back into a C++ one src/main.cpp an ordinary consumer, which imports the seam and @@ -34,6 +37,20 @@ one. The island also uses no standard library itself, which keeps it from linking a second copy of the C++ runtime into a program whose own copy came from mcpp's toolchain. +**That interface is generated, and the signatures exist once.** Each entry +point is marked with `MCPP_EXPORT_C` where it is defined; `mcpp.tools.island` +reads the marked declarations out of both implementations and writes the header +the island's compiler reads and the module the seam imports. A hand-written +header states each signature a second time, at the one boundary where a +disagreement is invisible: C language linkage does not mangle and the two halves +are never in one link, so two that disagreed would build cleanly and the +artifact would read its arguments by whichever signature it was compiled with. +The generator is handed both halves and refuses that there. + +`examples/09-heterogeneous/hip` keeps the hand-written header for the contrast. +Everything else about the two examples is the same computation, so the +difference between them is exactly this. + **The seam exists for backend substitution, not for the module boundary.** It is the one place where the island underneath becomes a CPU implementation, or could become HIP, without any consumer of `app.saxpy` changing, and the one @@ -74,7 +91,7 @@ The rule names it. This project writes one edge and no payload list at all: ```toml [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-cuda"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-cuda", "tools-island"], host-module = true } ``` `mcpp.rules.cuda` declares nvcc, cudart, cuRAND's headers, CCCL and the driver diff --git a/examples/09-heterogeneous/cuda/app/build.mcpp b/examples/09-heterogeneous/cuda/app/build.mcpp index 7eaf67d2c..1ff04deb5 100644 --- a/examples/09-heterogeneous/cuda/app/build.mcpp +++ b/examples/09-heterogeneous/cuda/app/build.mcpp @@ -1,19 +1,81 @@ import std; import mcpp; +import mcpp.tools.island; import mcpp.rules.cuda; +// THE BOUNDARY IS GENERATED, AND BOTH HALVES ARE READ. +// +// A seam has two implementations of one `extern "C"` boundary -- the island and +// the host fallback -- and exactly one of them is in any link. Both are handed +// to `scan` unconditionally: both exist on disk in either build, and which one +// is compiled is the manifest's decision, which this program does not repeat. +// +// What that buys is a check nothing else performs. C language linkage does not +// mangle and the two halves never meet at the link, so two that declared one +// name differently would build cleanly and the artifact would read its +// arguments by whichever signature it was compiled with. `scan` is the only +// point at which both texts exist at once, and it refuses there. +// +// The island reads the generated header through the compiler's forced-include +// flag, so no source in this project names a generated file and the project has +// no header in its source tree at all. +// WHERE THE FLAGS GO, AND WHY IT IS NOT `mcpp::cflag`. +// +// `force_include_flags` makes a compiler read the generated boundary header +// before the island's first line. The island is compiled by nvcc or by clang in +// CUDA mode -- a driver mcpp did not invoke -- so the flags have to reach THAT +// command line, which is what `mcpp::rules::cuda::options::flags` is for. +// +// The project-wide channels do not work here, and the reason is stronger than +// breadth: `mcpp::cxxflag` would force the header into every C++ translation +// unit, including `src/app.cppm`, and a module interface unit must begin with +// `export module`. Declarations ahead of that line are ill-formed. +static std::optional generate_boundary() { + const std::vector halves{ + std::string(mcpp::manifest_dir()) + "/src/kernels/saxpy.cu", + std::string(mcpp::manifest_dir()) + "/src/cpu/saxpy.cpp", + }; + for (auto const& f : halves) mcpp::rerun_if_changed(f.c_str()); + + mcpp::tools::island::options opt; + // Named rather than defaulted, so this project's modules share one root: + // `app.saxpy` is the seam and `app.kernels` is the boundary under it. The + // default derives from the package name, which would give + // `cuda_saxpy.kernels`. + opt.module_name = "app.kernels"; + opt.out_dir = std::string(mcpp::out_dir()) + "/island"; + opt.produced_by = "the cuda example"; + + const auto entries = mcpp::tools::island::scan(halves, opt); + if (!entries) return std::nullopt; + auto out = mcpp::tools::island::emit(*entries, opt); + if (!out) return std::nullopt; + + // The host half is ordinary project C++ and includes the generated header + // by name, so this directory has to be on the C++ include path too. + mcpp::include_dir(out->include_dir.c_str()); + mcpp::generated(out->interface_file.c_str()); + return out; +} + // Everything the rule needs is in the manifest: the architectures in -// `[build] accel`, the device sources in the constrained glob, the toolkit -// under `[xlings.workspace]`. This program names the island's include -// directory and says "go". +// `[build] accel`, the device sources in the constrained glob, the toolkit the +// rule itself declares. This program generates the boundary and says "go". // // The route follows the project's toolchain (clang for an LLVM toolchain, // nvcc for a GCC one); MCPP_EXAMPLE_CUDA_ROUTE=nvcc|clang overrides it, which // is how the alternate route is measured without editing the manifest. int main() { + const auto boundary = generate_boundary(); + if (!boundary) return 1; + mcpp::rerun_if_env_changed("MCPP_EXAMPLE_CUDA_ROUTE"); mcpp::rules::cuda::options opt; - opt.includes = { "include" }; + // NO `opt.includes`. The island's only header is the generated boundary, + // and it arrives through the forced-include flags below rather than through + // a search path -- the project has no `include/` tree left to name. + opt.flags = mcpp::tools::island::force_include_flags(boundary->header_file, + mcpp::compiler()); if (const char* r = std::getenv("MCPP_EXAMPLE_CUDA_ROUTE"); r && *r) opt.which = std::string_view(r) == "nvcc" ? mcpp::rules::cuda::route::nvcc : mcpp::rules::cuda::route::clang; diff --git a/examples/09-heterogeneous/cuda/app/include/saxpy/saxpy.h b/examples/09-heterogeneous/cuda/app/include/saxpy/saxpy.h deleted file mode 100644 index 447765eb2..000000000 --- a/examples/09-heterogeneous/cuda/app/include/saxpy/saxpy.h +++ /dev/null @@ -1,31 +0,0 @@ -// The device island's interface. -// -// `extern "C"` and free of standard-library types, on purpose. The island is -// compiled by nvcc driving a host compiler that mcpp did not choose, so the two -// sides do not share a C++ ABI and must not exchange anything that depends on -// one. Keeping the boundary this narrow is also what lets the island publish a -// C-surface compatibility tag. -#ifndef MCPP_EXAMPLE_SAXPY_H -#define MCPP_EXAMPLE_SAXPY_H - -#ifdef __cplusplus -extern "C" { -#endif - -// out[i] = a * x[i] + y[i], computed on the device. Returns 0 on success. -int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n); - -// WHICH DEVICE THE LAST SUCCESSFUL `saxpy_device` RAN ON, or "" if none has. -// -// A program that computes on a device and does not say where it computed -// cannot be checked. Both implementations of this seam produce the same four -// numbers, so the numbers alone do not distinguish a device run from the CPU -// fallback -- which is exactly the confusion an example about heterogeneous -// compute must not teach. Reading it before a successful call returns "", -// because a device run that did not happen has no device to name. -const char* saxpy_device_name(void); - -#ifdef __cplusplus -} -#endif -#endif diff --git a/examples/09-heterogeneous/cuda/app/mcpp.toml b/examples/09-heterogeneous/cuda/app/mcpp.toml index d1efb8304..162794045 100644 --- a/examples/09-heterogeneous/cuda/app/mcpp.toml +++ b/examples/09-heterogeneous/cuda/app/mcpp.toml @@ -19,8 +19,11 @@ default = "llvm@22.1.8" # The rule that compiles the island lives in the official plugin collection, # selected by its feature; `build.mcpp` imports it as `mcpp.rules.cuda`. +# `tools-island` is the second member this project uses: it writes the +# `extern "C"` boundary the island's compiler reads and the module the seam +# imports, so neither signature is stated twice. [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-cuda"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-cuda", "tools-island"], host-module = true } # The driver's userspace library, reached through an index package that owns # the one hop mcpp needs: a directory on the artifact's runtime search path. @@ -35,7 +38,7 @@ cuda-driver = "2026.09.05" # `cfg(accelerator = ...)` and the feature that selects it, so this project # names the rule and nothing else. To pin a different version, write the same # entry here and it wins -- see `One package, one version` in -# docs/05-mcpp-toml.md. Needs mcpp 2026.9.6.6 and mcpp:plugins 0.2.4. +# docs/05-mcpp-toml.md. Needs mcpp 2026.9.7.1 and mcpp:plugins 0.3.0. # # What the rule declares for this lane: nvcc, cudart, cuRAND's headers, CCCL and # the driver sentinel. The 12.9 line is an EXACT default rather than a floor, @@ -55,7 +58,8 @@ sources = [ "src/*.cpp", { glob = "src/kernels/**/*.cu", accel = "cuda12.9+{sm_89}" }, ] -include_dirs = ["include"] +# NO `include_dirs`. This project has no header of its own: the one the island +# reads is generated, and `build.mcpp` puts its directory on the include path. # The CUDA runtime is linked STATICALLY, and only when a device build asks for # it. Linking the redistributable half in leaves exactly one host dependency, diff --git a/examples/09-heterogeneous/cuda/app/src/app.cppm b/examples/09-heterogeneous/cuda/app/src/app.cppm index 5294eb3ae..6c5dd9977 100644 --- a/examples/09-heterogeneous/cuda/app/src/app.cppm +++ b/examples/09-heterogeneous/cuda/app/src/app.cppm @@ -5,10 +5,13 @@ // HIP or a CPU fallback without a single consumer of this module changing, and // a `cfg(accelerator = ...)` section has somewhere to apply. Remove the seam // and every importer becomes backend-specific. -module; -#include "saxpy/saxpy.h" export module app.saxpy; import std; +// The boundary, generated by `mcpp.tools.island` from the marked +// declarations in the two implementations. No global module fragment and +// no header: this seam imports the boundary the same way any consumer +// would, and each declaration exists once, beside its definition. +import app.kernels; export namespace app { diff --git a/examples/09-heterogeneous/cuda/app/src/cpu/saxpy.cpp b/examples/09-heterogeneous/cuda/app/src/cpu/saxpy.cpp index d56fef028..02c68239a 100644 --- a/examples/09-heterogeneous/cuda/app/src/cpu/saxpy.cpp +++ b/examples/09-heterogeneous/cuda/app/src/cpu/saxpy.cpp @@ -2,7 +2,20 @@ // asks for no accelerator (`mcpp build --no-accel`), through the // `cfg(not(accelerator = "cuda"))` section of the manifest; the device island // and this file define the same symbol and are never in one link. -#include "saxpy/saxpy.h" +// +// THIS FILE INCLUDES THE GENERATED BOUNDARY AND THE ISLAND DOES NOT, and the +// asymmetry is not an oversight. The island is compiled by a driver mcpp did +// not invoke, so its rule can be handed forced-include flags for that one +// command line. This half is ordinary project C++ compiled by mcpp, and the +// only project-wide channel available would force the header into every C++ +// translation unit -- including `src/app.cppm`, where declarations ahead of +// `export module` are ill-formed. One `#include` of a generated header in an +// ordinary source is the smaller thing. +// +// The header is generated from the marked declarations in BOTH files, so the +// signature below is still written once, and a signature here that disagreed +// with the island's is refused while the boundary is being generated. +#include "app.kernels.h" namespace { // Set by the call, read by the name. Empty until then, so `--no-accel` and a @@ -10,11 +23,13 @@ namespace { const char* g_ran_on = ""; } // namespace -extern "C" int saxpy_device(float a, const float* x, const float* y, - float* out, unsigned n) { +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]; g_ran_on = "cpu (this build names no accelerator)"; return 0; } -extern "C" const char* saxpy_device_name(void) { return g_ran_on; } +MCPP_EXPORT_C +const char* saxpy_device_name(void) { return g_ran_on; } diff --git a/examples/09-heterogeneous/cuda/app/src/kernels/saxpy.cu b/examples/09-heterogeneous/cuda/app/src/kernels/saxpy.cu index 4c9271f82..e765d3b6f 100644 --- a/examples/09-heterogeneous/cuda/app/src/kernels/saxpy.cu +++ b/examples/09-heterogeneous/cuda/app/src/kernels/saxpy.cu @@ -7,7 +7,6 @@ // second copy of the C++ runtime into a program whose own copy came from // mcpp's toolchain, which is the failure where one is linked and the other is // loaded. -#include "saxpy/saxpy.h" #include #include @@ -30,10 +29,12 @@ char g_ran_on[256] = ""; } // namespace -extern "C" const char* saxpy_device_name(void) { return g_ran_on; } +MCPP_EXPORT_C +const char* saxpy_device_name(void) { return g_ran_on; } -extern "C" int saxpy_device(float a, const float* x, const float* y, - float* out, unsigned n) { +MCPP_EXPORT_C +int saxpy_device(float a, const float* x, const float* y, + float* out, unsigned n) { float *dx = nullptr, *dy = nullptr, *dout = nullptr; const size_t bytes = static_cast(n) * sizeof(float); int rc = -1; diff --git a/examples/09-heterogeneous/hip/README.md b/examples/09-heterogeneous/hip/README.md index 96f725a44..bb5f8d53c 100644 --- a/examples/09-heterogeneous/hip/README.md +++ b/examples/09-heterogeneous/hip/README.md @@ -9,12 +9,39 @@ app/ src/kernels/saxpy.hip the island: a device translation unit written against the HIP API src/cpu/saxpy.cpp the same interface implemented for the host - include/saxpy/saxpy.h the island's interface: extern "C", no std types + include/saxpy/saxpy.h the island's interface: extern "C", no std types. + WRITTEN BY HAND, and that is why this example + exists beside `cuda/` -- see below src/app.cppm the seam src/main.cpp an ordinary consumer build.mcpp hands the device sources to `mcpp.rules.hip` ``` +## The boundary is written by hand here, and generated in `cuda/` + +The two examples compute the same thing through the same seam, so the one +difference between them is how the `extern "C"` boundary comes to exist. + +Here it is a file in the source tree. Each signature appears twice -- once in +`include/saxpy/saxpy.h` and once at each definition -- and nothing checks that +the copies agree. That is the ordinary arrangement, and it is worth seeing +written out, because it is what every C boundary looks like and because the +failure it permits is quiet: C language linkage does not mangle and the island +and the CPU fallback are never in one link, so a signature that drifted produces +a clean build and an artifact that reads its arguments by whichever version it +was compiled with. + +`cuda/` marks the entry points where they are defined and lets +`mcpp.tools.island` write the header and the module from them. The signatures +then exist once, and the generator -- handed both implementations -- refuses a +disagreement at the one point where both texts are in front of it. + +Neither is deprecated. A project whose boundary is stable, or whose island is +compiled somewhere mcpp cannot reach, writes the header; the generated form is +the default because the copy it removes is the one that goes wrong silently. + +## The kernel + Compare `src/kernels/saxpy.hip` with example 09's `saxpy.cu`: the same kernel, the same seam, and every device call spelled `hip*` instead of `cuda*`. diff --git a/examples/09-heterogeneous/hip/app/mcpp.toml b/examples/09-heterogeneous/hip/app/mcpp.toml index 4f19678ac..e982626eb 100644 --- a/examples/09-heterogeneous/hip/app/mcpp.toml +++ b/examples/09-heterogeneous/hip/app/mcpp.toml @@ -20,7 +20,7 @@ import_std = true default = "llvm@22.1.8" [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-hip"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-hip"], host-module = true } # The driver's userspace library. HIP reaches the device through the CUDA # runtime here, so this is the same one hop the CUDA consumer needs: mcpp's @@ -39,7 +39,7 @@ cuda-driver = "2026.09.05" # the rule's business rather than every consumer's. # # To pin a different version, write the same entry here and it wins. Needs mcpp -# 2026.9.6.6 and mcpp:plugins 0.2.4. +# 2026.9.7.1 and mcpp:plugins 0.3.0. [build] # Two chunks: the programming model, and the device. A device is spelled once diff --git a/examples/09-heterogeneous/multi-backend/README.md b/examples/09-heterogeneous/multi-backend/README.md index 892c0d4ac..34ea8bc1a 100644 --- a/examples/09-heterogeneous/multi-backend/README.md +++ b/examples/09-heterogeneous/multi-backend/README.md @@ -93,7 +93,7 @@ saying it after the vocabulary grows. ```toml [build-dependencies.mcpp] -plugins = { version = "0.2.2", features = ["rules-cuda", "rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-cuda", "rules-spirv"], host-module = true } ``` Two rules, in one build program, which is what an additive-backend package diff --git a/examples/09-heterogeneous/multi-backend/build.mcpp b/examples/09-heterogeneous/multi-backend/build.mcpp index f84901406..bba64ee14 100644 --- a/examples/09-heterogeneous/multi-backend/build.mcpp +++ b/examples/09-heterogeneous/multi-backend/build.mcpp @@ -1,5 +1,6 @@ import std; import mcpp; +import mcpp.plugins; import mcpp.rules.cuda; import mcpp.rules.spirv; @@ -26,5 +27,13 @@ int main() { mcpp::rules::spirv::options sp; sp.includes = { "src/backends/vulkan" }; + // THE PAYLOAD IS A MODULE. `src/backends/vulkan/host.cpp` opened with + // `#include "saxpy_comp.h"` -- a name no line in this project produced. + // It now imports `opkit.shaders`, named here rather than defaulted so this + // package's modules share the root its library module already uses; the + // default derives from the package name and would give + // `opkit_multi_backend.shaders`. + sp.surface = mcpp::plugins::surface::kind::module_; + sp.module_name = "opkit.shaders"; return mcpp::rules::spirv::compile(sp) ? 0 : 1; } diff --git a/examples/09-heterogeneous/multi-backend/mcpp.toml b/examples/09-heterogeneous/multi-backend/mcpp.toml index 9789570a5..401da1c37 100644 --- a/examples/09-heterogeneous/multi-backend/mcpp.toml +++ b/examples/09-heterogeneous/multi-backend/mcpp.toml @@ -40,7 +40,7 @@ default = "llvm@22.1.8" # The rules are declared unconditionally because `build.mcpp` imports them # unconditionally -- each returns immediately when its own backend is absent. [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-cuda", "rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-cuda", "rules-spirv"], host-module = true } # ── the payloads, gated on the device they are for ────────────────────────── # diff --git a/examples/09-heterogeneous/multi-backend/src/backends/vulkan/host.cpp b/examples/09-heterogeneous/multi-backend/src/backends/vulkan/host.cpp index 54931a5ce..7b1ed983e 100644 --- a/examples/09-heterogeneous/multi-backend/src/backends/vulkan/host.cpp +++ b/examples/09-heterogeneous/multi-backend/src/backends/vulkan/host.cpp @@ -20,8 +20,13 @@ // Generated by mcpp.rules.spirv into the build directory, which the rule puts on // the include path. The symbol name is derived from the file name: -// `src/backends/vulkan/saxpy.comp` -> `saxpy_comp_spv`. -#include "saxpy_comp.h" +// THE PAYLOAD ARRIVES AS A MODULE, NOT AS A FILE NAME. +// +// `opkit.shaders` is generated by `mcpp.rules.spirv` because `build.mcpp` asked +// for the module surface. `src/backends/vulkan/saxpy.comp` becomes +// `opkit::shaders::saxpy_comp()`, so this file names nothing the rule chose and +// the rule is free to change where it writes and how it stores the words. +import opkit.shaders; namespace { @@ -164,8 +169,12 @@ extern "C" int opkit_vulkan_saxpy(float a, const float* x, const float* y, { VkShaderModuleCreateInfo smci{}; smci.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO; - smci.codeSize = sizeof saxpy_comp_spv; - smci.pCode = saxpy_comp_spv; + // Address and size together, which is what makes `sizeof` the wrong + // question rather than an awkward one: under object storage there is no + // array to take the size of. + const auto spv = opkit::shaders::saxpy_comp(); + smci.codeSize = spv.size_bytes; + smci.pCode = spv.code; if (vkCreateShaderModule(dev, &smci, nullptr, &shader) != VK_SUCCESS) goto done; VkDescriptorSetLayoutBinding binding{}; diff --git a/examples/09-heterogeneous/sycl/README.md b/examples/09-heterogeneous/sycl/README.md index 81f7cffda..f36002259 100644 --- a/examples/09-heterogeneous/sycl/README.md +++ b/examples/09-heterogeneous/sycl/README.md @@ -10,7 +10,10 @@ app/ no BMI, and compiled by a SECOND compiler src/cpu/saxpy.cpp the same interface implemented for the host, compiled instead when the build asks for no accelerator - include/saxpy/saxpy.h the island's interface: extern "C", no std types + (generated) the `extern "C"` boundary and the module over it, + written by `mcpp.tools.island` from the marked + declarations in the two implementations above -- + there is no header in this source tree src/app.cppm the seam: a module that turns the C interface back into a C++ one src/main.cpp an ordinary consumer, which imports the seam and diff --git a/examples/09-heterogeneous/sycl/app/build.mcpp b/examples/09-heterogeneous/sycl/app/build.mcpp index 475d740d2..c4948a1f8 100644 --- a/examples/09-heterogeneous/sycl/app/build.mcpp +++ b/examples/09-heterogeneous/sycl/app/build.mcpp @@ -1,9 +1,61 @@ import std; import mcpp; +import mcpp.tools.island; import mcpp.rules.sycl; +// THE BOUNDARY IS GENERATED, AND BOTH HALVES ARE READ. +// +// A seam has two implementations of one `extern "C"` boundary -- the island and +// the host fallback -- and exactly one of them is in any link. Both are handed +// to `scan` unconditionally: both exist on disk in either build, and which one +// is compiled is the manifest's decision, which this program does not repeat. +// +// What that buys is a check nothing else performs. C language linkage does not +// mangle and the two halves never meet at the link, so two that declared one +// name differently would build cleanly and the artifact would read its +// arguments by whichever signature it was compiled with. `scan` is the only +// point at which both texts exist at once, and it refuses there. +static std::optional generate_boundary() { + const std::vector halves{ + std::string(mcpp::manifest_dir()) + "/src/kernels/saxpy.sycl", + std::string(mcpp::manifest_dir()) + "/src/cpu/saxpy.cpp", + }; + for (auto const& f : halves) mcpp::rerun_if_changed(f.c_str()); + + mcpp::tools::island::options opt; + // Named rather than defaulted, so this project's modules share one root: + // `app.saxpy` is the seam and `app.kernels` is the boundary under it. The + // default derives from the package name, which would give + // `sycl_saxpy.kernels` -- what `examples/09-heterogeneous/vulkan` uses for its + // shaders. + opt.module_name = "app.kernels"; + opt.out_dir = std::string(mcpp::out_dir()) + "/island"; + opt.produced_by = "the sycl example"; + + const auto entries = mcpp::tools::island::scan(halves, opt); + if (!entries) return std::nullopt; + auto out = mcpp::tools::island::emit(*entries, opt); + if (!out) return std::nullopt; + + // The host half is ordinary project C++ and includes the generated header + // by name, so this directory has to be on the C++ include path too. + mcpp::include_dir(out->include_dir.c_str()); + mcpp::generated(out->interface_file.c_str()); + return out; +} + int main() { + const auto boundary = generate_boundary(); + if (!boundary) return 1; + mcpp::rules::sycl::options opt; - opt.includes = { "include" }; + // NO `opt.includes`. The island's only header is the generated boundary, + // and it arrives through the forced-include flags below rather than through + // a search path -- the project has no `include/` tree left to name. + // + // NOT `mcpp::cflag`: the SYCL compiler is a driver mcpp did not invoke, and + // it inherits nothing from the project-wide flag channels. + opt.flags = mcpp::tools::island::force_include_flags(boundary->header_file, + mcpp::compiler()); return mcpp::rules::sycl::compile(opt) ? 0 : 1; } diff --git a/examples/09-heterogeneous/sycl/app/include/saxpy/saxpy.h b/examples/09-heterogeneous/sycl/app/include/saxpy/saxpy.h deleted file mode 100644 index 447765eb2..000000000 --- a/examples/09-heterogeneous/sycl/app/include/saxpy/saxpy.h +++ /dev/null @@ -1,31 +0,0 @@ -// The device island's interface. -// -// `extern "C"` and free of standard-library types, on purpose. The island is -// compiled by nvcc driving a host compiler that mcpp did not choose, so the two -// sides do not share a C++ ABI and must not exchange anything that depends on -// one. Keeping the boundary this narrow is also what lets the island publish a -// C-surface compatibility tag. -#ifndef MCPP_EXAMPLE_SAXPY_H -#define MCPP_EXAMPLE_SAXPY_H - -#ifdef __cplusplus -extern "C" { -#endif - -// out[i] = a * x[i] + y[i], computed on the device. Returns 0 on success. -int saxpy_device(float a, const float* x, const float* y, float* out, unsigned n); - -// WHICH DEVICE THE LAST SUCCESSFUL `saxpy_device` RAN ON, or "" if none has. -// -// A program that computes on a device and does not say where it computed -// cannot be checked. Both implementations of this seam produce the same four -// numbers, so the numbers alone do not distinguish a device run from the CPU -// fallback -- which is exactly the confusion an example about heterogeneous -// compute must not teach. Reading it before a successful call returns "", -// because a device run that did not happen has no device to name. -const char* saxpy_device_name(void); - -#ifdef __cplusplus -} -#endif -#endif diff --git a/examples/09-heterogeneous/sycl/app/mcpp.toml b/examples/09-heterogeneous/sycl/app/mcpp.toml index f075db6ca..b876124c9 100644 --- a/examples/09-heterogeneous/sycl/app/mcpp.toml +++ b/examples/09-heterogeneous/sycl/app/mcpp.toml @@ -20,7 +20,7 @@ import_std = true default = "llvm@22.1.8" [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-sycl"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-sycl", "tools-island"], host-module = true } # The SYCL runtime, on the artifact's runtime search path. mcpp's private # loader does not consult /usr/lib, so `libsycl.so.9` -- which the rule @@ -48,7 +48,7 @@ sycl-runtime = "2026.09.07" # resolves for exactly that version or for nothing. # # To pin a different version, write the same entry here and it wins. Needs mcpp -# 2026.9.6.6 and mcpp:plugins 0.2.4. +# 2026.9.7.1 and mcpp:plugins 0.3.0. [build] # Two chunks: the programming model, and the device. Written `sycl` alone, the @@ -60,7 +60,9 @@ sources = [ "src/*.cpp", { glob = "src/kernels/**/*.sycl", accel = "sycl, cuda12.9+{sm_89}" }, ] -include_dirs = ["include"] +# NO `include_dirs`. This project has no header of its own: the one the +# island reads is generated, and `build.mcpp` puts its directory on the +# include path. # `-lsycl` and the libstdc++ half are NOT here: they are consequences of the # compiler the rule chose, so the rule puts them on the link line itself diff --git a/examples/09-heterogeneous/sycl/app/src/app.cppm b/examples/09-heterogeneous/sycl/app/src/app.cppm index 137e6c422..1e8eeab22 100644 --- a/examples/09-heterogeneous/sycl/app/src/app.cppm +++ b/examples/09-heterogeneous/sycl/app/src/app.cppm @@ -5,10 +5,13 @@ // can become CUDA, SYCL or a CPU fallback without a single consumer of this // module changing, and a `cfg(accelerator = ...)` section has somewhere to // apply. Remove the seam and every importer becomes backend-specific. -module; -#include "saxpy/saxpy.h" export module app.saxpy; import std; +// The boundary, generated by `mcpp.tools.island` from the marked +// declarations in the two implementations. No global module fragment and +// no header: this seam imports the boundary the same way any consumer +// would, and each declaration exists once, beside its definition. +import app.kernels; export namespace app { diff --git a/examples/09-heterogeneous/sycl/app/src/cpu/saxpy.cpp b/examples/09-heterogeneous/sycl/app/src/cpu/saxpy.cpp index 41699168a..a61871e6e 100644 --- a/examples/09-heterogeneous/sycl/app/src/cpu/saxpy.cpp +++ b/examples/09-heterogeneous/sycl/app/src/cpu/saxpy.cpp @@ -2,7 +2,20 @@ // asks for no accelerator (`mcpp build --no-accel`), through the // `cfg(not(accelerator = "sycl"))` section of the manifest; the device island // and this file define the same symbol and are never in one link. -#include "saxpy/saxpy.h" +// +// THIS FILE INCLUDES THE GENERATED BOUNDARY AND THE ISLAND DOES NOT, and the +// asymmetry is not an oversight. The island is compiled by a driver mcpp did +// not invoke, so its rule can be handed forced-include flags for that one +// command line. This half is ordinary project C++ compiled by mcpp, and the +// only project-wide channel available would force the header into every C++ +// translation unit -- including `src/app.cppm`, where declarations ahead of +// `export module` are ill-formed. One `#include` of a generated header in an +// ordinary source is the smaller thing. +// +// The header is generated from the marked declarations in BOTH files, so the +// signature below is still written once, and one here that disagreed with the +// island's is refused while the boundary is being generated. +#include "app.kernels.h" namespace { // Set by the call, read by the name. Empty until then, so `--no-accel` and a @@ -10,11 +23,13 @@ namespace { const char* g_ran_on = ""; } // namespace -extern "C" int saxpy_device(float a, const float* x, const float* y, - float* out, unsigned n) { +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]; g_ran_on = "cpu (this build names no accelerator)"; return 0; } -extern "C" const char* saxpy_device_name(void) { return g_ran_on; } +MCPP_EXPORT_C +const char* saxpy_device_name(void) { return g_ran_on; } diff --git a/examples/09-heterogeneous/sycl/app/src/kernels/saxpy.sycl b/examples/09-heterogeneous/sycl/app/src/kernels/saxpy.sycl index d66a750c7..d61ce4003 100644 --- a/examples/09-heterogeneous/sycl/app/src/kernels/saxpy.sycl +++ b/examples/09-heterogeneous/sycl/app/src/kernels/saxpy.sycl @@ -29,7 +29,6 @@ // (`ProgramManager::getDeviceImage`) in neither of those two paths. That is // why the manifest names the device -- and why `mcpp.rules.sycl` warns at // build time when it does not. -#include "saxpy/saxpy.h" #include #include #include @@ -45,10 +44,12 @@ namespace { char g_ran_on[256] = ""; } // namespace -extern "C" const char* saxpy_device_name(void) { return g_ran_on; } +MCPP_EXPORT_C +const char* saxpy_device_name(void) { return g_ran_on; } -extern "C" int saxpy_device(float a, const float* x, const float* y, - float* out, unsigned n) { +MCPP_EXPORT_C +int saxpy_device(float a, const float* x, const float* y, + float* out, unsigned n) { // THE QUEUE TAKES AN ASYNC HANDLER, AND WITHOUT ONE THIS FUNCTION CANNOT // KEEP ITS PROMISE. // diff --git a/examples/09-heterogeneous/vulkan/README.md b/examples/09-heterogeneous/vulkan/README.md index 13c17cd24..1c66eae0a 100644 --- a/examples/09-heterogeneous/vulkan/README.md +++ b/examples/09-heterogeneous/vulkan/README.md @@ -14,7 +14,10 @@ app/ SPIR-V compiled into the binary as a C array src/cpu/saxpy.cpp the same interface implemented for the host, compiled instead when the build names no device - include/saxpy/saxpy.h the island's interface: extern "C", no std types + include/saxpy/saxpy.h the island's interface: extern "C", no std types. + Written by hand HERE, because this example is + about the other lane: the SPIR-V payload, which + arrives as a module and names no file at all src/app.cppm the seam: a module that turns the C interface back into a C++ one src/main.cpp an ordinary consumer, which imports the seam diff --git a/examples/09-heterogeneous/vulkan/app/build.mcpp b/examples/09-heterogeneous/vulkan/app/build.mcpp index 3ae0fdf81..3771690ea 100644 --- a/examples/09-heterogeneous/vulkan/app/build.mcpp +++ b/examples/09-heterogeneous/vulkan/app/build.mcpp @@ -1,7 +1,21 @@ import std; import mcpp; +import mcpp.plugins; import mcpp.rules.spirv; +// THE SURFACE IS A MODULE, AND THAT IS THE ONE LINE THAT DECIDES IT. +// +// `mcpp.rules.spirv` writes an interface the island imports instead of a header +// the island includes. The words in the artifact are identical; what changes is +// whether this project has to know the file name the rule chose. It did: +// `src/vulkan/saxpy.cpp` used to open with `#include "scale_comp.h"`, a name no +// line in this project produced and no reader could derive without opening the +// rule. +// +// The module name defaults to `.shaders`, so this +// package's shaders arrive as `vulkan_saxpy.shaders` and the accessor for +// `shaders/scale.comp` is `vulkan_saxpy::shaders::scale_comp()`. Both are +// derived from names this project already wrote. int main() { mcpp::rerun_if_changed_glob("shaders/**/*.comp"); mcpp::rerun_if_changed_glob("shaders/**/*.glsl"); @@ -9,5 +23,6 @@ int main() { mcpp::rules::spirv::options opt; // So a shader can `#include` another from the same directory. opt.includes = { "shaders" }; + opt.surface = mcpp::plugins::surface::kind::module_; return mcpp::rules::spirv::compile(opt) ? 0 : 1; } diff --git a/examples/09-heterogeneous/vulkan/app/mcpp.toml b/examples/09-heterogeneous/vulkan/app/mcpp.toml index f1fd7dfd7..820df29b8 100644 --- a/examples/09-heterogeneous/vulkan/app/mcpp.toml +++ b/examples/09-heterogeneous/vulkan/app/mcpp.toml @@ -13,7 +13,7 @@ import_std = true # The rule that compiles the shaders lives in the official plugin collection, # selected by its feature; `build.mcpp` imports it as `mcpp.rules.spirv`. [build-dependencies.mcpp] -plugins = { version = "0.2.4", features = ["rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-spirv"], host-module = true } # The Khronos loader, built by the index rather than taken from the host, and # the adapter that makes the host's own ICDs reachable from a binary running diff --git a/examples/09-heterogeneous/vulkan/app/src/vulkan/saxpy.cpp b/examples/09-heterogeneous/vulkan/app/src/vulkan/saxpy.cpp index d00910edc..5527d3014 100644 --- a/examples/09-heterogeneous/vulkan/app/src/vulkan/saxpy.cpp +++ b/examples/09-heterogeneous/vulkan/app/src/vulkan/saxpy.cpp @@ -3,11 +3,11 @@ // It is an island for the same reason the CUDA one is: the API underneath is // C, the memory it operates on is not the program's, and the seam above turns // both back into C++. What is different is where the device code came from — -// nothing here was compiled by the C++ toolchain. `scale_comp.h` is a -// `const uint32_t[]` that `mcpp.rules.spirv` produced from `shaders/scale.comp`, -// and the driver compiles it for whatever device is present when the program -// runs. That is the property this example exists to show: one artifact, no -// per-GPU build, and a CPU rasteriser is a legitimate device. +// nothing here was compiled by the C++ toolchain. The SPIR-V words come from +// `shaders/scale.comp` by way of `mcpp.rules.spirv`, and the driver compiles +// them for whatever device is present when the program runs. That is the +// property this example exists to show: one artifact, no per-GPU build, and a +// CPU rasteriser is a legitimate device. #include "saxpy/saxpy.h" #include @@ -17,10 +17,18 @@ #include #include -// Generated by mcpp.rules.spirv into the build directory, which the rule puts on -// the include path. The symbol name is derived from the file name: -// `shaders/scale.comp` -> `scale_comp_spv`. -#include "scale_comp.h" +// THE PAYLOAD ARRIVES AS A MODULE, NOT AS A FILE NAME. +// +// `vulkan_saxpy.shaders` is generated by `mcpp.rules.spirv` because +// `build.mcpp` asked for the module surface. The name is derived: the package +// is `vulkan-saxpy`, and `shaders/scale.comp` becomes `scale_comp()` under +// `vulkan_saxpy::shaders`. Nothing in this file names a generated path, so the +// rule is free to change where it writes and how it stores the words. +// +// An `import` in an ordinary `.cpp` is deliberate: this translation unit is not +// a module unit, it is the island's host half, and it reaches the payload the +// same way any consumer would. +import vulkan_saxpy.shaders; namespace { @@ -160,10 +168,16 @@ extern "C" int saxpy_device(float a, const float* x, const float* y, } { + // The accessor answers with both halves at once, which is what makes + // `sizeof` the wrong question rather than merely an awkward one: the + // rule may store the words in an object rather than in an array, and + // there is then no array to take the size of. + const auto spv = vulkan_saxpy::shaders::scale_comp(); + VkShaderModuleCreateInfo smci{}; smci.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO; - smci.codeSize = sizeof scale_comp_spv; - smci.pCode = scale_comp_spv; + smci.codeSize = spv.size_bytes; + smci.pCode = spv.code; if (vkCreateShaderModule(dev, &smci, nullptr, &shader) != VK_SUCCESS) goto done; VkDescriptorSetLayoutBinding binding{}; diff --git a/examples/10-graphics/README.md b/examples/10-graphics/README.md index 0423103f2..f3536f67b 100644 --- a/examples/10-graphics/README.md +++ b/examples/10-graphics/README.md @@ -53,7 +53,7 @@ tells the two apart, which is why it is printed. ```toml [build-dependencies.mcpp] -plugins = { version = "0.2.5", features = ["rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-spirv"], host-module = true } [build] accel = "vulkan1.2" @@ -65,10 +65,17 @@ sources = [ ``` That is the whole of it. `mcpp.rules.spirv` declares the shader compiler it -drives, so this project names no payload for it; the constrained globs route the -shaders to the build program rather than to the C++ compiler; and the generated -headers land on the include path, so `src/vulkan/render.cpp` writes -`#include "triangle_vert.h"`. +drives, so this project names no payload for it, and the constrained globs route +the shaders to the build program rather than to the C++ compiler. + +The compiled stages arrive as a MODULE. `build.mcpp` asks for that surface in +one line, and `src/vulkan/render.cpp` writes `import offscreen.shaders;` and +calls `offscreen::shaders::triangle_vert()`. It used to write +`#include "triangle_vert.h"` and `#include "triangle_frag.h"` -- two names no +line in this project produced and no reader could derive without opening the +rule. The accessor answers with the address and the byte count together, which +is what makes `sizeof` the wrong question rather than an awkward one: under +object storage there is no array to take the size of. **A dependency cannot be conditioned on the accelerator, and this project is where that shows.** `accelerator` is resolved from the dependency graph, so a diff --git a/examples/10-graphics/offscreen/build.mcpp b/examples/10-graphics/offscreen/build.mcpp index 31461b7f0..38a9857b2 100644 --- a/examples/10-graphics/offscreen/build.mcpp +++ b/examples/10-graphics/offscreen/build.mcpp @@ -1,13 +1,27 @@ import std; import mcpp; +import mcpp.plugins; import mcpp.rules.spirv; +// THE SHADERS ARE REACHED BY IMPORTING A MODULE. +// +// `src/vulkan/render.cpp` opened with `#include "triangle_vert.h"` and +// `#include "triangle_frag.h"` -- two names no line in this project produced, +// and which no reader could derive without opening the rule. One import +// replaces both, and the accessors are named after the files: +// `shaders/triangle.vert` becomes `offscreen::shaders::triangle_vert()`. +// +// `module_name` is stated rather than defaulted so the module root matches the +// project rather than its directory: the package is `offscreen-triangle`, and +// the default would be `offscreen_triangle.shaders`. int main() { mcpp::rerun_if_changed_glob("shaders/**/*.vert"); mcpp::rerun_if_changed_glob("shaders/**/*.frag"); mcpp::rerun_if_changed_glob("shaders/**/*.glsl"); mcpp::rules::spirv::options opt; - opt.includes = { "shaders" }; + opt.includes = { "shaders" }; + opt.surface = mcpp::plugins::surface::kind::module_; + opt.module_name = "offscreen.shaders"; return mcpp::rules::spirv::compile(opt) ? 0 : 1; } diff --git a/examples/10-graphics/offscreen/mcpp.toml b/examples/10-graphics/offscreen/mcpp.toml index 917fdd701..e8849d5b8 100644 --- a/examples/10-graphics/offscreen/mcpp.toml +++ b/examples/10-graphics/offscreen/mcpp.toml @@ -13,7 +13,7 @@ import_std = true # One edge. `mcpp.rules.spirv` declares the shader compiler it drives, so this # project names no payload for it. [build-dependencies.mcpp] -plugins = { version = "0.2.5", features = ["rules-spirv"], host-module = true } +plugins = { version = "0.3.0", features = ["rules-spirv"], host-module = true } # The Khronos loader. Unconditional, and the reason is a rule of the engine # rather than a preference: A DEPENDENCY CANNOT BE CONDITIONED ON A LAYER. @@ -62,19 +62,16 @@ vulkan-runtime = "2026.09.07" [target.'cfg(macos)'.xlings.workspace] "xim:moltenvk" = "1.4.2" -# WINDOWS DECLARES NO DEVICE, AND THAT IS MEASURED RATHER THAN AN OVERSIGHT. +# WINDOWS DECLARES NO DEVICE, AND WHAT RULED OUT THE OBVIOUS CAUSES IS IN +# `.github/workflows/ci-windows.yml` BESIDE THE STEP THAT WAS WITHDRAWN. # -# `xim:mesa-lavapipe` publishes a Windows payload and it installs correctly, so -# an entry here would work in the sense that the bytes arrive. The program was -# built and run with it declared, and printed `render unavailable` -- the report -# `src/main.cpp` makes when the render function returns nothing. +# Four things were measured and none of them is it: the loader is present (the +# program runs and reports for itself), the ICD manifest now parses (a real +# defect, fixed in openxlings/xim-pkgindex#781), `vulkan_lvp.dll` imports only +# system DLLs, and both spellings of the driver-files variable were set. # -# The failure is after the loader rather than at it. This program imports -# `vkCreateInstance` from `vulkan-1.dll`, so one that could not find that DLL -# would fail during image load and print nothing; it printed. What has not been -# established is why the payload's ICD enumerates no device under a process mcpp -# launched. Until it is, this platform builds the Vulkan half and runs the CPU -# fallback, which is what the `cfg(not(...))` section below carries. +# Declaring the driver here would download 56 MB that produces no device, so it +# is not declared until the remaining cause is known. [build] accel = "vulkan1.2" diff --git a/examples/10-graphics/offscreen/src/vulkan/render.cpp b/examples/10-graphics/offscreen/src/vulkan/render.cpp index a1efac9ea..da305e9ca 100644 --- a/examples/10-graphics/offscreen/src/vulkan/render.cpp +++ b/examples/10-graphics/offscreen/src/vulkan/render.cpp @@ -24,8 +24,14 @@ #include #include -#include "triangle_vert.h" -#include "triangle_frag.h" +// THE STAGES ARRIVE AS A MODULE, NOT AS TWO FILE NAMES. +// +// `offscreen.shaders` is generated by `mcpp.rules.spirv` because `build.mcpp` +// asked for the module surface. `shaders/triangle.vert` becomes +// `offscreen::shaders::triangle_vert()`, so nothing here names a path the rule +// chose -- and the rule stays free to change where it writes and whether the +// words live in a header or in an object. +import offscreen.shaders; namespace { @@ -269,8 +275,12 @@ extern "C" int render_offscreen(unsigned w, unsigned h, unsigned char* rgba) { if (vkCreateFramebuffer(f.device, &fbi, nullptr, &f.fb) != VK_SUCCESS) return 1; // ── the pipeline ──────────────────────────────────────────────────────── - f.vs = make_module(f.device, triangle_vert_spv, sizeof triangle_vert_spv); - f.fs = make_module(f.device, triangle_frag_spv, sizeof triangle_frag_spv); + // Address and size together. `sizeof` is not merely awkward here, it is + // unanswerable under object storage: there is no array to take the size of. + const auto vs = offscreen::shaders::triangle_vert(); + const auto fs = offscreen::shaders::triangle_frag(); + f.vs = make_module(f.device, vs.code, vs.size_bytes); + f.fs = make_module(f.device, fs.code, fs.size_bytes); if (!f.vs || !f.fs) return 1; VkPipelineShaderStageCreateInfo stages[2]{};