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<mcpp::tools::island::emitted> generate_boundary() {
+    const std::vector<std::string> 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 <cuda_runtime.h>
 #include <cstdio>
 
@@ -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<size_t>(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<mcpp::tools::island::emitted> generate_boundary() {
+    const std::vector<std::string> 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 <sycl/sycl.hpp>
 #include <cstdio>
 #include <string>
@@ -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 `<package with '-' as '_'>.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 <vulkan/vulkan.h>
@@ -17,10 +17,18 @@
 #include <cstring>
 #include <vector>
 
-// 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 <cstring>
 #include <vector>
 
-#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]{};