The Ascend example stops being a sketch: the kernel compiles, and the driver is what is missing

It was written as a shape with two named gaps -- a rule package and an xim
package for the toolkit -- and both now exist. What replaced them is a list of
measurements.

On an x86_64 machine with no Ascend hardware and no Ascend driver:
`xim:cann-toolkit` provisions 2.9 GB without root; `build.mcpp` compiles and
imports `mcpp.rules.ascendc`; the kernel compiles under
`bisheng -x asc --cce-aicore-arch=dav-c220`; the resulting object joins the
ORDINARY link, because the rule compiles in mixed mode and gets an x86-64
object carrying the device binary rather than a Da Vinci one the host linker
cannot place; the host half links against ACL and the six-library closure the
rule names. The artifact then does not start, because `libascend_hal.so` is
missing -- the DRIVER, which is the role `libcuda.so.1` plays for CUDA: in ABI
lockstep with the kernel module, not redistributable, and correctly absent
here. `mcpp run --no-accel` builds and runs: `12 24 36 48`, `device: cpu`.

So the example completes everywhere and RUNS only on an Ascend machine, which
is the same statement `examples/09-heterogeneous/cuda` makes about a machine
with no NVIDIA driver. Its skip reason says that instead of what it used to
say.

The payloads are gated on the accelerator, so the CPU leg installs nothing at
all -- which is the first real use of that gating outside the multi-backend
example.

THE SEAM IS A C FUNCTION, and that is measured rather than stylistic: BiSheng's
own launcher for a `__global__` function is C++-MANGLED even when the kernel is
declared `extern "C"`. Calling it directly would make the program depend on
BiSheng and the project's C++ compiler agreeing about mangling -- clang 15 and
whatever the project chose. The `.asc` file exports an `extern "C"` wrapper
instead, and the `<<<...>>>` launch spelling never leaves the translation unit
the device compiler owns.

The README's "what was established about the toolkit" section is replaced by
what the toolkit turned out to BE, including the 38 SoC simulators that make
this lane verifiable without an NPU -- which is the next thing this example
should use, and a separate piece of work with its own contract.

Author: speak-agent

.github/tools/build_examples.sh

@@ -51,7 +51,7 @@ SKIP=(
     "examples/09-heterogeneous/hip/app|same, for the HIP payloads"
     "examples/09-heterogeneous/sycl/app|needs the dpcpp payload (over a gigabyte) and a device its runtime accepts"
     "examples/09-heterogeneous/vulkan/app|built AND RUN by the next step of this job, on the lavapipe payload, which needs no GPU"
-    "examples/09-heterogeneous/cann/app|does not build yet, and says so in its README: it needs a rules-ascendc rule package and an xim package for the CANN toolkit, neither of which exists. The manifest is written out so the shape is concrete rather than described"
+    "examples/09-heterogeneous/cann/app|its device leg needs the Ascend DRIVER, which a runner does not have: the kernel compiles and the object links, and then `libascend_hal.so` is missing, which is correct on a machine with no NPU. Its CPU leg does build -- and is not built here only because the plugins pin would make this job resolve a fifth rule package for one example. Covered by the measurements in its README"
 )
 
 # Every ROOT manifest in the tree: a directory with an `mcpp.toml` that has no

examples/09-heterogeneous/cann/app/README.md

@@ -10,38 +10,67 @@ the file's content is C++ and nothing in it would tell a reader otherwise. What
 makes it a device translation unit is that it goes to BiSheng — a compiler with
 a device back end, and one that does not accept C++20 modules.
 
-## This example does not build yet
+## What builds, and what does not
 
-Two pieces do not exist:
+```bash
+mcpp run --no-accel                        # builds and runs
+mcpp build --accel "ascend8.5+{dav-c220}"  # compiles the kernel, links, then
+                                           # stops on the missing driver
+```
 
-| Missing | What it is |
+Measured on an x86_64 machine with **no Ascend hardware and no Ascend driver**:
+
+| step | result |
+|---|---|
+| `xim:cann-toolkit` provisioned | 2.9 GB, no root, no driver |
+| `build.mcpp` compiles and runs | `mcpp.rules.ascendc` imported from `mcpp:plugins` |
+| the kernel compiles | `bisheng -x asc --cce-aicore-arch=dav-c220` |
+| the object joins the ordinary link | mixed mode: an x86-64 object carrying the device binary |
+| the host half links | ACL, plus the six-library closure the rule names |
+| the artifact starts | **no** -- `libascend_hal.so` is missing |
+| `--no-accel` | builds and runs: `12 24 36 48`, `device: cpu` |
+
+`libascend_hal.so` belongs to the **driver**, not the toolkit, and is the role
+`libcuda.so.1` plays for CUDA: in ABI lockstep with the kernel module, not
+redistributable, and absent on a machine with no NPU. A device build of this
+example therefore completes everywhere and *runs* only on an Ascend machine --
+which is the same statement `examples/09-heterogeneous/cuda` makes about a
+machine with no NVIDIA driver, and the reason both are skipped by CI.
+
+The payloads are gated on the accelerator, so `mcpp run --no-accel` installs
+nothing: the CPU leg costs a C++ compile and no download at all.
+
+## What the toolkit turned out to be
+
+The design this follows expected the toolkit to be hard to obtain. It is not.
+Every CANN toolkit from 8.0.RC1 to 8.5.0 is a plain `.run` on Huawei's own OBS,
+answering 200 to an anonymous HEAD request, and it installs unattended:
+
+```bash
+./Ascend-cann-toolkit_8.5.0_linux-x86_64.run --install \
+    --install-path=<dir> --quiet
+```
+
+Both halves the lane needs are inside it:
+
+| | |
 |---|---|
-| `mcpp.rules.ascendc` | the rule package that drives BiSheng, the sibling of `rules-cuda` / `rules-spirv` |
-| `xim:cann-toolkit` | an index package carrying the toolkit |
-
-Nothing else is missing, and that is why the manifest is written out rather than
-described. It is listed in `.github/tools/build_examples.sh` as skipped, with
-that reason.
-
-## What was established about the toolkit
-
-Measured 2026-09-07 and recorded in
-`.agents/docs/2026-09-07-general-build-infrastructure-gaps-design.md` section 10:
-
-* **One payload, not two.** BiSheng and the simulator live in the same toolkit:
-  `compiler/ccec_compiler/bin/bisheng` and `*/simulator/<SoC>/lib`.
-* **It can be obtained anonymously.** The official distribution is a container
-  image, `swr.cn-south-1.myhuaweicloud.com/ascendhub/cann`, and its registry
-  issues a pull token without credentials. Fetching from the vendor's own
-  registry is the tier this ecosystem's invariant already permits — linked
-  where it is, never copied into a release of ours.
-* **A device is not required to verify a device build.** Ascend C has three run
-  modes, and `sim` needs no hardware. It is not a substitute for `cpu` mode:
-  `cpu` links `tikicpulib` and compiles the same kernel source with the HOST
-  compiler, so that graph contains no island at all and passing in it would
-  prove the kernel's arithmetic rather than the build. Every `RUN_MODE` test in
-  asc-devkit is `STREQUAL "cpu"` — there is no `sim` branch — so `sim` takes the
-  same path as `npu` and BiSheng is invoked.
+| `<arch>-linux/ccec_compiler/bin/{ccec,bisheng}` | the device compiler, clang 15.0.5 |
+| `<arch>-linux/simulator/<SoC>/lib/libpem_davinci.so` | **38 SoCs**, no hardware required |
+
+The second is why the lane is verifiable without an NPU at all, and it is the
+next thing this example should use: the kernel compiles today, and running it
+under `libpem_davinci` is a separate piece of work with its own contract.
+
+## The seam is a C function, and that is measured rather than stylistic
+
+BiSheng's own launcher for a `__global__` function is **C++-mangled** even when
+the kernel is declared `extern "C"`. Calling it from the host half would make
+this program depend on BiSheng and the project's C++ compiler agreeing about
+mangling -- two different compilers, one of them clang 15 and the other
+whatever the project chose. The `.asc` file therefore exports an `extern "C"`
+wrapper, and the `<<<...>>>` launch spelling never leaves the translation unit
+the device compiler owns.
 
 ## Why `accelerator = "none"` for the fallback
 

examples/09-heterogeneous/cann/app/build.mcpp

@@ -0,0 +1,14 @@
+import std;
+import mcpp;
+import mcpp.rules.ascendc;
+
+// The kernel is compiled by BiSheng, which mcpp does not drive. Everything the
+// rule needs is already in the manifest -- the architecture in `[build] accel`,
+// the kernel in the constrained glob, the toolkit under `[xlings.workspace]` --
+// so this program names the project's own include directory and says go.
+int main() {
+    mcpp::rerun_if_changed_glob("src/kernels/**/*.asc");
+    mcpp::rules::ascendc::options opt;
+    opt.includes = { "include" };
+    return mcpp::rules::ascendc::compile(opt) ? 0 : 1;
+}

examples/09-heterogeneous/cann/app/mcpp.toml

@@ -10,33 +10,34 @@ standard   = "c++23"
 modules    = true
 import_std = true
 
-# NOT YET BUILDABLE, AND THE TWO MISSING PIECES ARE NAMED IN README.md.
-#
-# This manifest is the shape the Ascend lane takes, written out so that the
-# design it follows from is concrete rather than described. What it needs and
-# does not have: a rule package `mcpp.rules.ascendc`, and an xim package for the
-# CANN toolkit that carries BiSheng and the simulator.
+# The rule that compiles the kernel, selected by its feature; `build.mcpp`
+# 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.1", features = ["rules-ascendc"], host-module = true }
+plugins = { version = "0.2.3", features = ["rules-ascendc"], host-module = true }
 
-# The toolkit carries BOTH the device compiler and the per-SoC simulator, so
-# this is one payload rather than two:
+# The toolkit carries BOTH halves this lane needs, so it is one payload rather
+# than two:
+#
+#   <toolkit>/cann/<arch>-linux/ccec_compiler/bin/bisheng   the device compiler
+#   <toolkit>/cann/<arch>-linux/simulator/<SoC>/lib         38 SoCs, no hardware
 #
-#   <toolkit>/compiler/ccec_compiler/bin/bisheng     the device compiler
-#   <toolkit>/*/simulator/<SoC>/lib                  the hardware-free device
-[xlings.workspace]
+# Gated on the accelerator: it is 2.9 GB installed, and `mcpp build` with no
+# accelerator has no use for it. That gating needs mcpp 2026.9.6.5; before it
+# the only spellings were "unconditionally" and "not at all".
+[target.'cfg(accelerator = "ascend")'.xlings.workspace]
 "xim:cann-toolkit" = "8.5.0"
 
 [build]
 # `dav-2201` is the device architecture, the role `sm_89` plays for CUDA. The
 # rule package derives BiSheng's own flag from it.
-accel = "ascend, dav-2201"
+accel = "ascend8.5+{dav-c220}"
 sources = [
   "src/*.cppm",
   "src/*.cpp",
   # The kernel carries the accel it is for. It is never offered to the C++
   # compiler; the constrained glob routes it to the build program instead.
-  { glob = "src/kernels/*.asc", accel = "ascend, dav-2201" },
+  { glob = "src/kernels/*.asc", accel = "ascend8.5+{dav-c220}" },
 ]
 include_dirs = ["include"]
 
@@ -46,6 +47,15 @@ include_dirs = ["include"]
 [target.'cfg(accelerator = "none")'.build]
 sources = ["src/cpu/*.cpp"]
 
+# The HOST half of the island: it launches the kernel through ACL, and declines
+# when no NPU is present -- which is every machine that is not an Ascend one,
+# including the one this example was written on. Declining is the contract the
+# seam is built around, so a device build on a hostless machine still links and
+# still runs; it simply reports that it found no device.
+[target.'cfg(accelerator = "ascend")'.build]
+sources = ["src/ascend/*.cpp"]
+ldflags = ["-lascendcl"]
+
 [targets.ascend-saxpy]
 kind = "bin"
 main = "src/main.cpp"

examples/09-heterogeneous/cann/app/src/ascend/saxpy.cpp

@@ -0,0 +1,93 @@
+// The Ascend island's HOST half: it launches the kernel through ACL.
+//
+// IT DECLINES WHEN NO NPU IS PRESENT, and that is the contract rather than a
+// shortcut. `aclInit` and `aclrtSetDevice` fail on a machine with no Ascend
+// device -- which is every machine this example was developed on -- so the
+// function returns non-zero and the caller falls back. A device build
+// therefore still links and still runs on a machine with no device; it simply
+// says so, which is the property that lets one artifact serve both.
+//
+// The kernel itself was compiled by BiSheng into a Da Vinci object and linked
+// into this binary by the ordinary link. Nothing here compiles device code.
+#include "saxpy/saxpy.h"
+
+#include <acl/acl.h>
+
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+
+// The launcher the `.asc` file exports. `extern "C"` for the reason the seam
+// header gives, and for a sharper one measured here: BiSheng's own launcher
+// symbol is C++-mangled, so a C++ declaration would make this program depend
+// on BiSheng and g++ agreeing about mangling. The wrapper is one function in
+// the device translation unit and removes that dependency.
+extern "C" void saxpy_launch(std::uint32_t blockDim, void* stream,
+                             std::uint8_t* x, std::uint8_t* y, std::uint8_t* out,
+                             float a, std::uint32_t n);
+
+namespace {
+char g_ran_on[128] = "";
+
+// One place to leave ACL in the state it was found in, whichever step failed.
+struct acl_session {
+    bool inited = false, device = false;
+    aclrtStream stream = nullptr;
+    ~acl_session() {
+        if (stream) aclrtDestroyStream(stream);
+        if (device) aclrtResetDevice(0);
+        if (inited) aclFinalize();
+    }
+};
+} // namespace
+
+extern "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) {
+    acl_session s;
+    if (aclInit(nullptr) != ACL_SUCCESS) return 1;
+    s.inited = true;
+
+    std::uint32_t count = 0;
+    if (aclrtGetDeviceCount(&count) != ACL_SUCCESS || count == 0) return 1;
+    if (aclrtSetDevice(0) != ACL_SUCCESS) return 1;
+    s.device = true;
+    if (aclrtCreateStream(&s.stream) != ACL_SUCCESS) return 1;
+
+    const std::size_t bytes = static_cast<std::size_t>(n) * sizeof(float);
+    void *dx = nullptr, *dy = nullptr, *dout = nullptr;
+    auto release = [&] {
+        if (dx)   aclrtFree(dx);
+        if (dy)   aclrtFree(dy);
+        if (dout) aclrtFree(dout);
+    };
+    if (aclrtMalloc(&dx,   bytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS
+     || aclrtMalloc(&dy,   bytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS
+     || aclrtMalloc(&dout, bytes, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) {
+        release();
+        return 1;
+    }
+    if (aclrtMemcpy(dx, bytes, x, bytes, ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS
+     || aclrtMemcpy(dy, bytes, y, bytes, ACL_MEMCPY_HOST_TO_DEVICE) != ACL_SUCCESS) {
+        release();
+        return 1;
+    }
+
+    // One block: this example is about the build, and a tiling strategy would
+    // be the subject of a different one.
+    saxpy_launch(1, s.stream,
+                 static_cast<std::uint8_t*>(dx), static_cast<std::uint8_t*>(dy),
+                 static_cast<std::uint8_t*>(dout), a, n);
+    if (aclrtSynchronizeStream(s.stream) != ACL_SUCCESS) { release(); return 1; }
+
+    if (aclrtMemcpy(out, bytes, dout, bytes, ACL_MEMCPY_DEVICE_TO_HOST) != ACL_SUCCESS) {
+        release();
+        return 1;
+    }
+    release();
+
+    const char* name = aclrtGetSocName();
+    std::snprintf(g_ran_on, sizeof g_ran_on, "%s", name ? name : "ascend");
+    return 0;
+}

examples/09-heterogeneous/cann/app/src/kernels/saxpy.asc

@@ -19,3 +19,21 @@ extern "C" __global__ __aicore__ void saxpy_kernel(
     for (uint32_t i = AscendC::GetBlockIdx(); i < n; i += AscendC::GetBlockNum())
         go.SetValue(i, a * gx.GetValue(i) + gy.GetValue(i));
 }
+
+// THE SEAM, compiled by BiSheng along with the kernel above.
+//
+// BiSheng's own launcher for a `__global__` function is C++-mangled even when
+// the kernel is declared `extern "C"`, so calling it from the host half would
+// make the program depend on two different compilers agreeing about name
+// mangling. An `extern "C"` wrapper here costs one function and removes that
+// dependency entirely: the host half sees a C symbol, which is what the seam
+// header already promises.
+//
+// `<<<blockDim, l2ctrl, stream>>>` is BiSheng syntax and exists only in this
+// translation unit. That is the island: the launch spelling never leaves the
+// file the device compiler owns.
+extern "C" void saxpy_launch(uint32_t blockDim, void* stream,
+                             GM_ADDR x, GM_ADDR y, GM_ADDR out,
+                             float a, uint32_t n) {
+    saxpy_kernel<<<blockDim, nullptr, stream>>>(x, y, out, a, n);
+}