diff --git a/.agents/docs/2026-09-05-heterogeneous-build-ecosystem-design-v2.md b/.agents/docs/2026-09-05-heterogeneous-build-ecosystem-design-v2.md index 3a1cd394..9ab71fad 100644 --- a/.agents/docs/2026-09-05-heterogeneous-build-ecosystem-design-v2.md +++ b/.agents/docs/2026-09-05-heterogeneous-build-ecosystem-design-v2.md @@ -451,8 +451,8 @@ Status is one of `done`, `open`, `deferred (reason)`. | N2 | `examples/11-sycl-compute`: one SYCL kernel behind a seam, the CUDA backend and the host device from one artifact | `12 24 36 48` on the RTX 4080 and again with `--no-accel` | N1, Q1 | open | | N3 | `examples/12-hip-compute`: a `.hip` kernel through `mcpp.rules.hip` on the NVIDIA platform | `12 24 36 48`, and no `/usr` path on any command line | Q2, Z2 | open | | N4 | Chapter 20 gains the SYCL and HIP lanes and the table of which rule drives which compiler; both languages | `check_docs_style.sh` passes; the Chinese chapter has the rows the English one has | - | done: the lane table, the two-chunk `accel` explanation, the HIP header-layer note and the two-C++-runtimes note, in both languages; "Not implemented" corrected -- the whole-target shape is no longer among the missing | -| N5 | Version 2026.9.6.1, CHANGELOG, unit and e2e suites, PR, CI green, self-review | no e2e failure other than the ones already failing on `origin/main` | N1-N4 | open | -| N6 | Release, mirror, index bump, bootstrap pin | both mirrors return 200 and identical bytes for every asset; xim-pkgindex names it `latest` | N5 | open | +| N5 | Version 2026.9.6.1, CHANGELOG, unit and e2e suites, PR, CI green, self-review | no e2e failure other than the ones already failing on `origin/main` | N1-N4 | done: PR #574, 36/36 checks green, merged as `12d4759`; main green on every workflow at that sha except `ci-fresh-install`, which fails in `apt-get` inside a debian:11 container on five consecutive commits including two that predate this round | +| N6 | Release, mirror, index bump, bootstrap pin | both mirrors return 200 and identical bytes for every asset; xim-pkgindex names it `latest` | N5 | done: all four artefacts byte-identical GitHub/GitCode by GET; bump PR #770 merged (`6c2f283`) with `Publish Index Artifact` green; `xlings install mcpp@2026.9.6.1` yields `mcpp 2026.9.6.1` at the store path; pin on main | #### mcpp-plugins (single PR, version 0.2.0) @@ -461,22 +461,22 @@ Status is one of `done`, `open`, `deferred (reason)`. | Q1 | `mcpp.rules.sycl`, feature `rules-sycl` | a fixture compiles a `.sycl` unit with the `dpcpp` payload and links it into a C++23-modules program; the rule refuses with the declaration to add when the payload is absent | N1 | done: `12 24 36 48` on an RTX 4080 through `mcpp run`, zero `/usr` paths on any command line, and the device-link wrapper asserted as its own file | | Q2 | `mcpp.rules.hip`, feature `rules-hip` | a fixture compiles a `.hip` unit on the NVIDIA platform; the rule names the platform it selected and why | Z2 | done: `12 24 36 48` on an RTX 4080, zero `/usr` paths; the AMD platform is refused by name with the reason rather than approximated | | Q3 | `mcpp.rules.spirv` gains the `glslc` route now that a payload exists | both compilers produce a header the same program includes; the rule states which one it used | Z1 | done: the two command lines share almost nothing -- glslc's `-mfmt=c` is a bare initialiser list and its `-S` means "emit assembly" where glslang's names a stage -- so the rule writes the declaration itself and keys its `mcpp::fact` on the flavour | -| Q4 | CI: one consumer fixture per feature, built with the pinned mcpp | green on the PR head | Q1-Q3, N6 | open | -| Q5 | Release `v0.2.0`; GitHub archive and a GitCode asset with identical bytes; index descriptor | both URLs return 200 and one sha256 | Q4 | open | +| Q4 | CI: one consumer fixture per feature, built with the pinned mcpp | green on the PR head | Q1-Q3, N6 | done, and it found three host leaks the previous check could not see -- see 7.9 | +| Q5 | Release `v0.2.0`; GitHub archive and a GitCode asset with identical bytes; index descriptor | both URLs return 200 and one sha256 | Q4 | done: PR #3 merged (`56d9da2`), tag v0.2.0, both mirrors 47,792 bytes under one sha256 (`b5ee0cf4`) | #### mcpp-index (single PR) | # | task | criterion | depends on | status | |---|---|---|---|---| | R1 | `compat.sycl-runtime` 2026.09.06; then `mcpp:plugins` 0.2.0 in `pkgs/m/mcpp.plugins.lua`, floor 2026.9.6.1 | the adapter: a workspace member loads `libsycl.so.9` by soname under mcpp's own loader. the descriptor: `mcpp add mcpp:plugins` resolves in a sandbox and the new features are selectable | Q5, N6 | adapter done (PR #354 merged as `e17fc88`, index artifact published); descriptor open | -| R2 | CI green; merge; index artifact published | `Publish Index Artifact` green on the merge commit | R1 | open | +| R2 | CI green; merge; index artifact published | `Publish Index Artifact` green on the merge commit | R1 | done: PR #356 merged (`e68c9d8`), artifact green. The snapshot needed the whole `data/` directory removed to move -- clearing the version marker and the cache JSON was not enough, and neither mirror was at fault | #### Verification | # | task | criterion | depends on | status | |---|---|---|---|---| -| V3 | A fresh subos, CN mirror for both tools: install mcpp 2026.9.6.1 from the index, build examples 09 through 12 against `mcpp:plugins` 0.2.0 | every example builds; example 10 answers on lavapipe; the device half of 09, 11 and 12 compiles and says cleanly that the sandbox has no device | Z4, N6, Q5, R2 | open | -| V4 | The same on this host with the GPU | `12 24 36 48` from 09, 11 and 12 on the RTX 4080 | V3 | open | +| V3 | A fresh subos, CN mirror for both tools: install mcpp 2026.9.6.1 from the index, build examples 11 and 12 against `mcpp:plugins` 0.2.0 | every example builds; the device half compiles and the sandbox says cleanly that it has no device | Z4, N6, Q5, R2 | **done, 0 assertions failed.** Both examples built and answered `12 24 36 48` under `--no-accel`; the device-link wrapper was produced; no `/usr` on any command line. It also found two defects of mine -- see 7.9 | +| V4 | The same on this host with the GPU | `12 24 36 48` from 11 and 12 on the RTX 4080 | V3 | **done, 0 assertions failed.** Both examples on the device, `sycl-ls` enumerating the RTX 4080 from the store, and the adapter farm complete down to `libcuda.so.1` | ### 7.4 The completion rule @@ -587,3 +587,167 @@ mcpp-plugins, then mcpp-index (descriptor), then mcpp (examples). | 2026-09-06 | the round is five PRs, not three | the release order is a cycle otherwise | | 2026-09-06 | dpcpp corrected a second time (xim #769) | the first correction fixed the payload's own programs and broke every consumer of it; the regression was invisible from the payload | | 2026-09-06 | `compat.sycl-runtime` 2026.09.07 (mcpp-index #355) | a SYCL project should not have to declare CUDA; the adapter owns the hop its runtime needs | +| 2026-09-06 | the SYCL island gains an async handler and an inner catch; the rule warns when `accel` names sycl and no device | exercising the rule's other branch showed the island could not keep the promise its own comment made; see 7.7 | + +### 7.7 What exercising the second branch found + +`mcpp.rules.sycl` has two branches: `accel = "sycl, cuda12.9+{sm_89}"` compiles +ahead of time for that architecture, and `accel = "sycl"` compiles to SPIR-V and +leaves the device to the runtime. The examples take the first. Taking the second +built correctly and then died on an RTX 4080 with + + terminate called after throwing an instance of 'sycl::_V1::exception' + terminate called recursively + +and no message of the program's own -- contradicting the fixture's own comment, +which said the island turns its failures into a return code. + +Two hypotheses were wrong before the backtrace settled it, and both were worth +fixing anyway because each terminates on some machine: + +* An exception unwinding through the three `sycl::buffer` destructors. A buffer + destructor blocks until the work that reads it has finished, so a catch + outside the scope unwinds through three destructors that are themselves + finishing failed work. The catch moved inside the scope. +* The queue's asynchronous handler. SYCL delivers most device errors to it, and + a queue constructed without one gets the default handler, which calls + `std::terminate`. No catch intercepts that, because it never travels as an + exception through the frame. A handler is installed. + +The actual cause is neither. `ProgramManager::getDeviceImage` throws inside the +SYCL scheduler because the CUDA back end does not consume SPIR-V, and that path +reaches neither the caller's frame nor the queue's handler. Nothing the island +can write catches it. + +So the rule says it at build time instead, as an advisory rather than a +refusal: the SPIR-V form is correct and portable, and the rule cannot know the +machine's devices -- which is exactly why the statement belongs before the +first kernel rather than in a crash with no message. + +The general shape, which is the reusable part: **a promise a comment makes +about error handling is only tested by the configuration that fails.** Every +run on this machine took the path where nothing goes wrong, and the comment +read as true for as long as that held. + +### 7.8 Ecosystem review: what a consumer sees after this round + +A project that wants a GPU now writes the same four things whichever +programming model it picks, and the fourth is the only one that names the +model. + +```toml +[dependencies.mcpp] +plugins = { version = "0.2.0", features = ["rules-sycl"], host-module = true } + +[dependencies.compat] +sycl-runtime = "2026.09.07" # the runtime the artifact reaches through the loader + +[xlings.workspace] # payloads: the project picks the versions +"xim:dpcpp" = "7.1.0" +"xim:gcc" = "15.1.0" +"xim:cuda-nvcc" = "12.9.86" + +[build] +accel = "sycl, cuda12.9+{sm_89}" +``` + +Four properties hold across all four lanes, and each is enforced by something +rather than asserted here. + +**The device is spelled once.** `cuda12.9+{sm_89}` is character for character +the same in the CUDA, HIP and SYCL manifests. The first chunk names the +programming model; the second names the device. A rule reads the chunk it owns +and ignores the rest, so adding a fifth model adds no spelling. + +**The engine still knows no vendor.** `test_core_vendor_probes` strips comments +from `src/` and refuses any vendor tool name, with the file count as its own +denominator. Three rules and four payloads later, that test is unchanged. + +**No device command line reaches the host.** Asserted per lane in the plugin +collection's CI as a separate step from the build, because all three lanes +compiled successfully on a developer machine while reading `/usr/include` -- +twice in this round alone, once for the SYCL unit's C++ standard library and +once for clang's CUDA installation. Measured here: zero `/usr` paths across +CUDA, HIP and SYCL. + +**A payload is reachable, not merely installed.** The round's largest single +defect was a payload whose five reporting programs could not start; the round's +second was the correction to it, which fixed the payload and broke every +consumer. Both are now assertions: the programs start and carry DT_RPATH, and +the payload's libraries carry no search path at all. + +Three things this round could not close, each with what would decide it: + +| open | what would decide it | +|---|---| +| Intel `anv` in the Mesa payload | a libclc and SPIRV-LLVM-Translator chain, then a Mesa rebuild; both packages exist on conda-forge and were confirmed repackable, and are withheld because nothing yet consumes them | +| the AMD platform of HIP | a ROCm runtime and device library in `xim-pkgindex`. The rule refuses it by name today rather than emitting an object nothing can link | +| llama.cpp's Vulkan lane | its blocker is gone -- `glslc` is published -- and the remaining work is in that project's own `-m` repository rather than here | + +And one thing measured rather than assumed: `accel = "sycl"` without a device +compiles to SPIR-V, which the CUDA back end does not consume. The build is +correct, the failure is the machine's, and the rule now says so before the +first kernel. See 7.7. + +### 7.9 What CI and the two verification runs found, and the shape they share + +Eight defects in this round, six of them in work written for it, and not one +was found by reading the code. The list is worth keeping because the *method* +that found each is the reusable part. + +| found by | defect | +|---|---| +| using the payload | `xim:dpcpp` shipped five programs that could not start | +| CI's dependency-closure check | the first repair swapped the interpreter, losing the host fallback | +| rebuilding the example | the second repair fixed the payload and broke every consumer of it | +| reading `clang -###` | `-lstdc++` is rewritten to `-lc++` and vanishes | +| a runner | `libz.so.1` missing from the farm; this machine had it | +| taking the rule's other branch | the SYCL island's error promise held only where nothing failed | +| a runner | the HIP device pass read the host's `/usr/include/c++/14` | +| the sandbox run | example 11 carried a superseded manifest shape, and section G measured a farm the example had never used | + +THE SHAPE MOST OF THEM SHARE is that a criterion pointed at the wrong object, +and the wrong object answered `ok`: + +* the host-leak check read the command line, and an implicit include search is + never on it -- so it reported clean on the machine that was leaking; +* the farm check picked a directory out of the store by sorting, and a store + that has seen two adapter versions holds two farms; +* a payload lookup named one namespace, and the payload was installed under the + other; +* a program that could not start had its loader error attributed to an adapter, + turning one defect into three reports of which the third named the wrong + cause. + +Round 3's notes already record this shape twice ("the report was chosen by +`ls | head -1`", "the mirror was read from the tool's stdout"). It recurred +four times here, which suggests the lesson is not "watch for this" but +something mechanical: **when a check reports on an object it selected, it must +print which object it selected.** Section G now does, and the line +`note: falling back to the newest farm in the store` is what turned the last of +these from a false pass into a visible one. + +A property the same run exposed in a merged package: `compat.sycl-runtime` has +one `install()` and never reads `pkginfo.version()`, so a version key selects +the anchor URL and nothing else. The 2026.09.07 commit said 2026.09.06 +"installs the farm it always did"; it does not, and mcpp-index #357 corrects +the comment rather than the behaviour, which is intended for an adapter whose +content is generated. + +### 7.10 One failure this round did not cause and did not fix + +`ci-fresh-install` is red on mcpp's main and has been for five consecutive +commits, two of them before this round began. It fails in `apt-get install` +inside a `debian:11` container, in a step named "Install prerequisites" that +runs before mcpp is involved: + + E: Failed to fetch .../libperl5.32_5.32.1-4+deb11u5_amd64.deb 404 Not + Found [IP: 146.75.38.132] + +Measured rather than assumed: that exact URL returns 200 from here, and +`bullseye-security/InRelease` is still served -- so the package exists and a +CDN node was answering 404. The step already runs `apt-get update`, so the +mitigation would be `-o Acquire::Retries=3` rather than a fresh index. It is +not applied here: Docker Hub is unreachable from this machine, the failure does +not reproduce from this network, and an unverifiable change to an unrelated +workflow is worth less than a precise diagnosis of it. diff --git a/.agents/docs/2026-09-06-round4-verify.sh b/.agents/docs/2026-09-06-round4-verify.sh index b94a1b17..4dd17dfd 100755 --- a/.agents/docs/2026-09-06-round4-verify.sh +++ b/.agents/docs/2026-09-06-round4-verify.sh @@ -64,7 +64,7 @@ section "B. xim:shaderc and xim:hip-nvidia" # glslc: the criterion is a SPIR-V module's magic number, not an exit status -- # a compiler that wrote an empty file would also exit 0. "$XL" install shaderc -y >/dev/null 2>&1 || true -glslc=$(ls "$HOME"/.xlings/data/xpkgs/xim-x-shaderc/*/bin/glslc "$HOME"/.mcpp/registry/data/xpkgs/xim-x-shaderc/*/bin/glslc 2>/dev/null | head -1) +glslc=$(ls "$HOME"/.xlings/data/xpkgs/*-x-shaderc/*/bin/glslc "$HOME"/.mcpp/registry/data/xpkgs/*-x-shaderc/*/bin/glslc 2>/dev/null | head -1) if [ -n "$glslc" ]; then ok "xim:shaderc installed at $glslc" printf '#version 450\nlayout(local_size_x=64) in;\nlayout(std430,binding=0) buffer B { uint v[]; };\nvoid main(){ v[gl_GlobalInvocationID.x] *= 2; }\n' > "$work/s.comp" @@ -76,7 +76,15 @@ if [ -n "$glslc" ]; then fi # Every DT_NEEDED inside a payload. This is the property that makes the # repack a payload rather than a copy of somebody's /usr. - outside=$(ldd "$glslc" 2>/dev/null | awk '/=>/ {print $3}' | grep -v '^$' | grep -v "$HOME" | grep -v '^(' | head -3) + # THE LOADER IS NOT A DEPENDENCY RESOLVED FROM THE HOST. glibc's `ldd` + # prints the program interpreter with `=>` like everything else, so a + # filter that only looks for that arrow counts `/lib64/ld-linux-x86-64.so.2` + # as a host library and fails a payload that is in fact complete. The + # interpreter is chosen by the ELF header, not searched for, and which one + # a payload uses is the subject of its own recipe rather than of this + # check. + outside=$(ldd "$glslc" 2>/dev/null | awk '/=>/ {print $3}' | grep -v '^$' \ + | grep -v "$HOME" | grep -v '^(' | grep -v 'ld-linux' | head -3) [ -z "$outside" ] && ok "every glslc dependency resolves inside a payload" || fail "glslc resolves outside the store: $(echo "$outside" | tr '\n' ' ')" else fail "xim:shaderc left no bin/glslc in either store" @@ -85,7 +93,7 @@ fi # hip-nvidia: headers only, and the dispatch header is what makes the NVIDIA # platform reachable at all. "$XL" install hip-nvidia -y >/dev/null 2>&1 || true -hipinc=$(ls -d "$HOME"/.xlings/data/xpkgs/xim-x-hip-nvidia/*/include "$HOME"/.mcpp/registry/data/xpkgs/xim-x-hip-nvidia/*/include 2>/dev/null | head -1) +hipinc=$(ls -d "$HOME"/.xlings/data/xpkgs/*-x-hip-nvidia/*/include "$HOME"/.mcpp/registry/data/xpkgs/*-x-hip-nvidia/*/include 2>/dev/null | head -1) if [ -n "$hipinc" ]; then ok "xim:hip-nvidia installed at $hipinc" for h in hip/hip_runtime.h hip/hip_version.h hip/nvidia_detail/nvidia_hip_runtime.h hip/amd_detail/amd_hip_runtime_pt_api.h; do @@ -106,7 +114,7 @@ fi # devices are listed is the machine's answer and is deliberately not asserted. section "C. xim:dpcpp" "$XL" install dpcpp -y >/dev/null 2>&1 || true -dp=$(ls -d "$HOME"/.xlings/data/xpkgs/xim-x-dpcpp/*/ "$HOME"/.mcpp/registry/data/xpkgs/xim-x-dpcpp/*/ 2>/dev/null | head -1) +dp=$(ls -d "$HOME"/.xlings/data/xpkgs/*-x-dpcpp/*/ "$HOME"/.mcpp/registry/data/xpkgs/*-x-dpcpp/*/ 2>/dev/null | head -1) if [ -n "$dp" ]; then ok "xim:dpcpp installed at $dp" for p in sycl-ls sycl-prof sycl-trace sycl-sanitize syclbin-dump; do @@ -116,11 +124,34 @@ if [ -n "$dp" ]; then *"error while loading shared libraries"*) fail "bin/$p cannot start: $err" ;; *) ok " bin/$p starts" ;; esac + # DT_RPATH, not DT_RUNPATH, and the difference decides the adapters: + # RUNPATH is honoured for the program's own DT_NEEDED and not for a + # dlopen beneath it, so a RUNPATH here lets every one of these start + # and report no devices at all. + tag=$(readelf -d "$dp/bin/$p" 2>/dev/null | grep -oE 'RPATH|RUNPATH' | head -1) + [ "$tag" = "RPATH" ] && ok " bin/$p carries DT_RPATH" || fail "bin/$p carries '$tag', not RPATH" done + # And the LIBRARIES must have none: a RUNPATH on one of them switches off + # the inherited RPATH of whatever loaded it, which is what cut a consumer's + # artifact off from its own driver farm. + adapter=$(ls "$dp"/lib/libur_adapter_*.so.0.* 2>/dev/null | head -1) + if [ -n "$adapter" ]; then + n=$(readelf -d "$adapter" 2>/dev/null | grep -cE 'RPATH|RUNPATH') + [ "$n" -eq 0 ] && ok " the adapters carry no search path of their own, so they inherit the caller's" \ + || fail "$(basename "$adapter") carries a search path; it would cut its loader off from the artifact's farm" + fi + # ONE DEFECT, ONE REPORT, AT ITS CAUSE. A program that could not start + # leaves `error while loading shared libraries: libsycl.so.9` on the same + # stderr the adapter check reads, so running both turns one failure into + # three and the third names the wrong thing. lsout=$("$dp/bin/sycl-ls" --verbose 2>&1) - for own in libsycl.so libumf.so libur_loader.so; do - printf '%s' "$lsout" | grep -q "$own.*cannot open" && fail "an adapter could not find $own, which is in this payload" || ok " no adapter failed on $own" - done + if printf '%s' "$lsout" | grep -q 'error while loading shared libraries'; then + printf 'note: adapter checks skipped -- sycl-ls itself could not start (reported above)\n' + else + for own in libsycl.so libumf.so libur_loader.so; do + printf '%s' "$lsout" | grep -q "$own.*cannot open" && fail "an adapter could not find $own, which is in this payload" || ok " no adapter failed on $own" + done + fi if [ -n "${MCPP_VERIFY_HOST:-}" ]; then printf '%s' "$lsout" | grep -q 'cuda:gpu' && ok "sycl-ls enumerates a CUDA device on this host" || fail "sycl-ls found no CUDA device on a host that has one" else @@ -197,7 +228,7 @@ fi # -- F. the SYCL lane --------------------------------------------------------- section "F. examples/11-sycl-kernel" if [ -n "$SRC" ] && [ -d "$SRC/examples/11-sycl-kernel/app" ]; then - ex="$work/ex11"; cp -r "$SRC/examples/11-sycl-kernel/app" "$ex" + ex="$work/ex11"; cp -r "$SRC/examples/11-sycl-kernel/app" "$ex"; ex11_built="$ex" out=$(cd "$ex" && "$STORE" build --no-accel 2>&1 && "$STORE" run --no-accel 2>&1) printf '%s\n' "$out" | grep -q '12 24 36 48' && ok "example 11 --no-accel answered 12 24 36 48" \ || fail "example 11 --no-accel: $(printf '%s' "$out" | tail -3 | tr '\n' ' ')" @@ -234,14 +265,34 @@ fi # runtime loads and enumerates nothing), `libstdc++.so.6` (which compat.cudart # deliberately does not farm, and which a SYCL artifact needs because it links # libc++) and `libz.so.1` (which a developer machine happened to have). -section "G. compat.sycl-runtime" -farm=$(ls -d "$HOME"/.mcpp/registry/data/xpkgs/compat-x-sycl-runtime/*/mcpp_generated/sycl_runtime/lib 2>/dev/null | sort -V | tail -1) -if [ -z "$farm" ] && [ -n "$SRC" ]; then - farm=$(ls -d "$SRC"/examples/11-sycl-kernel/app/.mcpp/.xlings/data/xpkgs/compat-x-sycl-runtime/*/mcpp_generated/sycl_runtime/lib 2>/dev/null | sort -V | tail -1) +section "G. compat.sycl-runtime 2026.09.07" +# THE FARM THE EXAMPLE RESOLVED, NOT WHATEVER THE STORE HOLDS. +# +# A store that has seen two adapter versions holds two farms, and picking one +# by sorting is picking by accident: the sandbox run of 2026-09-06 reported +# `ok` for every soname of a farm the example under test had not used, because +# a DIFFERENT version happened to be the only one in that store. The example's +# own `resolution.json` names the directory it was actually built against. +farm="" +if [ -n "${ex11_built:-}" ]; then + farm=$(python3 - "$ex11_built" <<'PY' 2>/dev/null +import glob, json, sys +for f in glob.glob(sys.argv[1] + "/target/*/*/resolution.json"): + for m in json.dumps(json.load(open(f))).split('"'): + if m.endswith("sycl_runtime/lib"): + print(m); raise SystemExit +PY +) +fi +# Only if the example was not built here: then any installed farm is the best +# available evidence, and the line says which one it read. +if [ -z "$farm" ]; then + farm=$(ls -d "$HOME"/.mcpp/registry/data/xpkgs/compat-x-sycl-runtime/*/mcpp_generated/sycl_runtime/lib 2>/dev/null | sort -V | tail -1) + [ -n "$farm" ] && printf 'note: no example build to read; falling back to the newest farm in the store\n' fi if [ -n "$farm" ]; then ok "farm at $farm" - for so in libsycl.so.9 libur_loader.so.0 libumf.so.1 libstdc++.so.6 libz.so.1 libdl.so.2; do + for so in libsycl.so.9 libur_loader.so.0 libumf.so.1 libstdc++.so.6 libz.so.1 libdl.so.2 libcuda.so.1; do [ -e "$farm/$so" ] && ok " $so" || fail "the farm is missing $so" done # No unversioned name: mcpp puts runtime.library_dirs on the LINK line as diff --git a/docs/01-examples.md b/docs/01-examples.md index 1935929f..481cc0d0 100644 --- a/docs/01-examples.md +++ b/docs/01-examples.md @@ -29,6 +29,8 @@ examples. | 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-cuda-kernel`](../examples/09-cuda-kernel/) | A CUDA kernel behind a seam module, with a CPU fallback | `accel`, constrained source globs, `mcpp::action` with `role = "object"`, `cfg(accelerator = …)` | | 10 | [`examples/10-vulkan-compute`](../examples/10-vulkan-compute/) | The same computation as a Vulkan compute shader, on a GPU or on the CPU | `mcpp.rules.spirv` from `mcpp:plugins`, `mcpp::action` with `role = "source"`, generated headers, a software driver as a payload | +| 11 | [`examples/11-sycl-kernel`](../examples/11-sycl-kernel/) | The same computation as a SYCL kernel, compiled by a second compiler | `mcpp.rules.sycl` from `mcpp:plugins`, the `.sycl` device extension, a chained `mcpp::action` for the device link, `compat:sycl-runtime` | +| 12 | [`examples/12-hip-kernel`](../examples/12-hip-kernel/) | The same computation in HIP, reaching an NVIDIA device | `mcpp.rules.hip` from `mcpp:plugins`, HIP as a header layer over the CUDA runtime, a two-chunk `accel` | ## Suggested Reading Order diff --git a/docs/zh/01-examples.md b/docs/zh/01-examples.md index 86f37ea8..7beffc4c 100644 --- a/docs/zh/01-examples.md +++ b/docs/zh/01-examples.md @@ -26,6 +26,8 @@ mcpp build && mcpp run | 08 | [`examples/08-build-rules`](../../examples/08-build-rules/) | 两个规则包,以及同时用到它们的工程 | `host-module = true`、`[build-dependencies]`、`role = "check"` 的 `mcpp::action` | | 09 | [`examples/09-cuda-kernel`](../../examples/09-cuda-kernel/) | 接缝模块背后的 CUDA kernel,并带 CPU 回退 | `accel`、带约束的 source glob、`role = "object"` 的 `mcpp::action`、`cfg(accelerator = …)` | | 10 | [`examples/10-vulkan-compute`](../../examples/10-vulkan-compute/) | 同一个计算写成 Vulkan compute shader,在 GPU 上或在 CPU 上 | 来自 `mcpp:plugins` 的 `mcpp.rules.spirv`、`role = "source"` 的 `mcpp::action`、生成的头文件、作为载荷的软件驱动 | +| 11 | [`examples/11-sycl-kernel`](../../examples/11-sycl-kernel/) | 同一个计算写成 SYCL kernel,由第二个编译器编译 | 来自 `mcpp:plugins` 的 `mcpp.rules.sycl`、`.sycl` 设备扩展名、为 device link 串起来的 `mcpp::action`、`compat:sycl-runtime` | +| 12 | [`examples/12-hip-kernel`](../../examples/12-hip-kernel/) | 同一个计算写成 HIP,够到一台 NVIDIA 设备 | 来自 `mcpp:plugins` 的 `mcpp.rules.hip`、HIP 作为 CUDA 运行时之上的一层头文件、两段式的 `accel` | ## 推荐阅读顺序 diff --git a/examples/11-sycl-kernel/README.md b/examples/11-sycl-kernel/README.md new file mode 100644 index 00000000..6b98b6ff --- /dev/null +++ b/examples/11-sycl-kernel/README.md @@ -0,0 +1,102 @@ +# 11 — A SYCL kernel behind the same seam + +What this example demonstrates, and what it does not. + +## The shape + +``` +app/ + src/kernels/saxpy.sycl the island: a device translation unit, never scanned, + 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 + 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 + never sees the header + build.mcpp hands the device sources to `mcpp.rules.sycl`, a + member of `mcpp:plugins` selected by `rules-sycl` +``` + +Everything above is the shape of example 09, one file name apart. That is the +point of this example: SYCL is a different compilation model and it reaches +the build through the same seam, the same constrained glob and the same rule +mechanism. + +## What makes a `.sycl` file a device unit + +Not its content. Open `src/kernels/saxpy.sycl` and it is ordinary C++ — there +is no dialect to see, no `__global__`, no launch syntax. What makes it a device +translation unit is that it goes to a compiler with a device back end, which +mcpp does not drive and which does not accept C++20 modules. `SourceKind::Device` +states exactly that property and nothing about the language. + +Naming it `.cpp` and routing it by glob was possible and was rejected: one +extension would then mean two things depending on which glob matched first, and +the seam is legible precisely because the file name says which side of it a +unit is on. + +## Two edges, not one per source + +A SYCL object carries its device image, and nothing registers that image with +the runtime. The registration comes from a **device link** (`-fsycl-link`), +which reads every device object and emits one further host object. So the rule +submits one action per source and one that consumes their outputs, and the +engine orders them by the graph rather than by declaration order. + +Without the second, this program links, starts, and finds no kernel. + +## Three payloads, and what each one closes + +```toml +[xlings.workspace] +"xim:dpcpp" = "7.1.0" # the compiler: its clang has the SYCL front end +"xim:gcc" = "15.1.0" # the C++ standard library the unit compiles against +"xim:cuda-nvcc" = "12.9.86" # the NVIDIA back end's libdevice +``` + +`xim:gcc` is not a second toolchain. Left alone, the SYCL compiler takes its +C++ standard library headers from the host's GCC, and `--cuda-path` aside, it +finds the host's CUDA installation the same way. Neither says anything when it +happens: both are visible only in the compiler's own include search list, and +only on a machine that has those directories. That is why the rule refuses +without them and names the line to add. + +## Two C++ runtimes, and why the seam is not optional here + +`libsycl.so` is compiled against libstdc++ while an mcpp artifact links libc++, +so both are in the image. mcpp's duplicate-symbol check reports the unwinder +symbols they share, and the warning is correct. + +Nothing may cross the seam. The island catches its own `sycl::exception` and +returns a code, because the runtime that threw it is not the one the caller +would unwind with. Example 09's island can promise not to touch the standard +library at all; this one cannot — SYCL *is* a C++ library — so the discipline +moves from "no standard library" to "nothing crosses". + +Making that promise true took three things, and only two of them are a `catch`: + +* the catch sits **inside** the buffer scope, because a `sycl::buffer` + destructor blocks until the work that reads it has finished, and unwinding + through three of those is a second throw during unwinding; +* the queue takes an **asynchronous handler**, because a queue constructed + without one gets the default handler, and the default handler calls + `std::terminate` — which no `catch` can intercept, since it never travels as + an exception through this frame; +* and one failure remains outside both. A build compiled to SPIR-V, run against + a back end that does not consume it, throws from inside the SYCL scheduler. + That is why this manifest names the device, and why `mcpp.rules.sycl` warns at + build time when an `accel` names `sycl` and no device. + +## Running it + +``` +mcpp build # ahead of time for sm_89, through the dpcpp payload +mcpp run # 12 24 36 48, on the device +mcpp build --no-accel # the constrained glob is left out +mcpp run --no-accel # 12 24 36 48, from src/cpu/saxpy.cpp +``` + +A machine with no device runs the second pair and gets the same answer. That is +what the seam buys, and it is the same sentence example 09's README ends with. diff --git a/examples/11-sycl-kernel/app/build.mcpp b/examples/11-sycl-kernel/app/build.mcpp new file mode 100644 index 00000000..475d740d --- /dev/null +++ b/examples/11-sycl-kernel/app/build.mcpp @@ -0,0 +1,9 @@ +import std; +import mcpp; +import mcpp.rules.sycl; + +int main() { + mcpp::rules::sycl::options opt; + opt.includes = { "include" }; + return mcpp::rules::sycl::compile(opt) ? 0 : 1; +} diff --git a/examples/11-sycl-kernel/app/include/saxpy/saxpy.h b/examples/11-sycl-kernel/app/include/saxpy/saxpy.h new file mode 100644 index 00000000..8d82d972 --- /dev/null +++ b/examples/11-sycl-kernel/app/include/saxpy/saxpy.h @@ -0,0 +1,21 @@ +// 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); + +#ifdef __cplusplus +} +#endif +#endif diff --git a/examples/11-sycl-kernel/app/mcpp.toml b/examples/11-sycl-kernel/app/mcpp.toml new file mode 100644 index 00000000..67b5700a --- /dev/null +++ b/examples/11-sycl-kernel/app/mcpp.toml @@ -0,0 +1,71 @@ +[package] +name = "sycl-saxpy" +namespace = "example" +version = "0.1.0" +description = "A SYCL kernel behind a seam module, reaching an NVIDIA device, with a CPU fallback" +accelerators = ["sycl"] + +[language] +standard = "c++23" +modules = true +import_std = true + +# The C++ half is mcpp's own clang. The device half cannot be: SYCL needs a +# compiler with the SYCL front end, which is what the dpcpp payload is, and +# that second compiler is the entire reason this rule exists. +[toolchain] +default = "llvm@22.1.8" + +[dependencies.mcpp] +plugins = { version = "0.2.0", features = ["rules-sycl"], 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 +# satisfied at LINK time from the payload -- would not be found at run time. +# +# ONE compat entry, and that is the point of where the boundary was drawn. The +# SYCL runtime's CUDA back end dlopens the driver, so `compat:sycl-runtime` +# carries that hop itself from 2026.09.07 on. A project that writes SYCL does +# not have to know CUDA is underneath it. +[dependencies.compat] +sycl-runtime = "2026.09.07" + +# Three payloads, each closing one hole the host would otherwise fill. +# +# dpcpp the compiler. +# gcc NOT a second toolchain: the C++ standard library the SYCL unit +# compiles against. Left alone, dpcpp's clang reads the HOST's +# /usr/include/c++ -- measured, and invisible until a machine +# without it refuses the build. +# cuda-nvcc the NVIDIA back end's libdevice. Also measured: without +# `--cuda-path` clang finds the host's CUDA installation and says +# nothing about it. +[xlings.workspace] +"xim:dpcpp" = "7.1.0" +"xim:gcc" = "15.1.0" +"xim:cuda-nvcc" = "12.9.86" + +[build] +# Two chunks: the programming model, and the device. Written `sycl` alone, the +# unit is compiled to SPIR-V and the runtime picks a device at run time; with +# the second chunk it is compiled ahead of time for that architecture. +accel = "sycl, cuda12.9+{sm_89}" +sources = [ + "src/*.cppm", + "src/*.cpp", + { glob = "src/kernels/**/*.sycl", accel = "sycl, cuda12.9+{sm_89}" }, +] +include_dirs = ["include"] + +# `-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 +# through `mcpp::link_lib`. A manifest that named them would be stating a fact +# about a compiler it did not select. + +# The CPU-only variant: the same seam, a host implementation behind it. +[target.'cfg(not(accelerator = "sycl"))'.build] +sources = ["src/cpu/*.cpp"] + +[targets.sycl-saxpy] +kind = "bin" +main = "src/main.cpp" diff --git a/examples/11-sycl-kernel/app/src/app.cppm b/examples/11-sycl-kernel/app/src/app.cppm new file mode 100644 index 00000000..a5e1a3ff --- /dev/null +++ b/examples/11-sycl-kernel/app/src/app.cppm @@ -0,0 +1,27 @@ +// The seam. +// +// Its reason for existing is not that a device compiler rejects modules. It is +// that this is the one place a backend can be exchanged: the island underneath +// 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; + +export namespace app { + +// The device interface is raw pointers and a count because it has to be. The +// seam is where that becomes a C++ interface again. +std::optional> +saxpy(float a, std::span x, std::span y) { + if (x.size() != y.size()) return std::nullopt; + std::vector out(x.size()); + if (saxpy_device(a, x.data(), y.data(), out.data(), + static_cast(x.size())) != 0) + return std::nullopt; + return out; +} + +} // namespace app diff --git a/examples/11-sycl-kernel/app/src/cpu/saxpy.cpp b/examples/11-sycl-kernel/app/src/cpu/saxpy.cpp new file mode 100644 index 00000000..1881cbf1 --- /dev/null +++ b/examples/11-sycl-kernel/app/src/cpu/saxpy.cpp @@ -0,0 +1,11 @@ +// The CPU implementation behind the same seam. Compiled only when the build +// 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" + +extern "C" int saxpy_device(float a, const float* x, const float* y, + float* out, unsigned n) { + for (unsigned i = 0; i < n; ++i) out[i] = a * x[i] + y[i]; + return 0; +} diff --git a/examples/11-sycl-kernel/app/src/kernels/saxpy.sycl b/examples/11-sycl-kernel/app/src/kernels/saxpy.sycl new file mode 100644 index 00000000..2f8f1fe4 --- /dev/null +++ b/examples/11-sycl-kernel/app/src/kernels/saxpy.sycl @@ -0,0 +1,101 @@ +// The island, written in SYCL. Nothing here is visible to the module graph: a +// device translation unit is never scanned and never produces a BMI. +// +// THIS ISLAND USES THE C++ STANDARD LIBRARY, AND ITS SIBLINGS DELIBERATELY DO +// NOT. +// +// tests/cuda-consumer and tests/hip-consumer say in their own comments that +// their islands touch no standard library, because an island that pulls one in +// links a second C++ runtime into a program whose own came from mcpp's +// toolchain. A SYCL island cannot make that promise: SYCL IS a C++ library, +// `sycl::queue` and `sycl::buffer` are class templates over standard types, +// and `libsycl.so` is itself compiled against libstdc++. +// +// So the second runtime is present here by construction, and what makes that +// safe is the seam rather than an absence. Nothing below crosses +// `saxpy_device`'s signature: the parameters are pointers and a count, the +// return is an int, and no object allocated by one runtime is freed by the +// other. `mcpp.rules.sycl` links the libstdc++ half explicitly for the same +// reason -- it is a consequence of the compiler this rule chose, not something +// the manifest should have to know. +// +// WHAT THIS ISLAND CAN AND CANNOT TURN INTO A RETURN CODE. +// +// It catches the synchronous half (no usable device, a rejected submit) and +// installs a handler for the asynchronous half, which SYCL otherwise delivers +// to a default handler that calls `std::terminate`. What it cannot catch is a +// missing device image: a build compiled to SPIR-V, run against a back end +// that does not consume SPIR-V, throws from inside the scheduler +// (`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 + +extern "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. + // + // SYCL splits errors in two. A synchronous one -- no usable device, a bad + // submit -- is thrown where it happens and an ordinary catch sees it. An + // ASYNCHRONOUS one, which is most of what a device reports, is delivered + // to the queue's exception handler; a queue constructed WITHOUT one gets + // the default handler, and the default handler calls `std::terminate`. + // + // It is not something a catch can intercept, because it does not travel as + // an exception through this frame. Measured, on a build whose SPIR-V no + // device on this machine could consume, with every `try` below already in + // place: + // + // terminate called after throwing an instance of 'sycl::_V1::exception' + // terminate called recursively + // + // A device island that promises the rest of the program a return code has + // to install a handler, or the promise is only kept on machines where + // nothing goes wrong. + int rc = 0; + auto on_async = [&rc](sycl::exception_list errors) { + for (std::exception_ptr const& e : errors) { + try { std::rethrow_exception(e); } + catch (const sycl::exception& ex) { + std::fprintf(stderr, "sycl (async): %s\n", ex.what()); + } + } + rc = -1; + }; + + try { + sycl::queue q{on_async}; + { + sycl::buffer bx(x, sycl::range<1>(n)); + sycl::buffer by(y, sycl::range<1>(n)); + sycl::buffer bo(out, sycl::range<1>(n)); + // The catch is INSIDE the buffer scope on purpose: a buffer + // destructor blocks until the work that reads it has finished, and + // letting an exception propagate through three of those means + // unwinding through destructors that are themselves finishing + // failed work. Caught here, the scope ends normally. + try { + q.submit([&](sycl::handler& h) { + sycl::accessor ax(bx, h, sycl::read_only); + sycl::accessor ay(by, h, sycl::read_only); + sycl::accessor ao(bo, h, sycl::write_only, sycl::no_init); + h.parallel_for(sycl::range<1>(n), + [=](sycl::id<1> i) { ao[i] = a * ax[i] + ay[i]; }); + }); + q.wait_and_throw(); + } catch (const sycl::exception& e) { + std::fprintf(stderr, "sycl: %s\n", e.what()); + rc = -1; + } + } + } catch (const sycl::exception& e) { + // The queue itself: a machine with no device the runtime can use fails + // here, before any buffer exists. + std::fprintf(stderr, "sycl: no usable device: %s\n", e.what()); + rc = -1; + } + return rc; +} diff --git a/examples/11-sycl-kernel/app/src/main.cpp b/examples/11-sycl-kernel/app/src/main.cpp new file mode 100644 index 00000000..6efc1938 --- /dev/null +++ b/examples/11-sycl-kernel/app/src/main.cpp @@ -0,0 +1,12 @@ +import std; +import app.saxpy; + +int main() { + const std::vector x{1, 2, 3, 4}, y{10, 20, 30, 40}; + auto out = app::saxpy(2.0f, x, y); + if (!out) { std::println("device unavailable"); return 1; } + for (auto v : *out) std::print("{} ", v); + std::println(""); + const std::vector want{12, 24, 36, 48}; + return *out == want ? 0 : 1; +} diff --git a/examples/12-hip-kernel/README.md b/examples/12-hip-kernel/README.md new file mode 100644 index 00000000..50b749ce --- /dev/null +++ b/examples/12-hip-kernel/README.md @@ -0,0 +1,76 @@ +# 12 — A HIP kernel on the NVIDIA platform + +What this example demonstrates, and what it does not. + +## The shape + +``` +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 + src/app.cppm the seam + src/main.cpp an ordinary consumer + build.mcpp hands the device sources to `mcpp.rules.hip` +``` + +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*`. + +## HIP on this platform is a header layer, not a second runtime + +HIP has two implementations behind one API. On AMD hardware it is a runtime +library that talks to ROCm. On NVIDIA hardware every entry point is an inline +wrapper over the CUDA one — `hipMalloc` resolves to `cudaMalloc` through the +header, `hipError_t` is `cudaError_t` under a typedef — so the object links +against the CUDA runtime and nothing of ROCm's. + +Three consequences, all visible in this example's manifest: + +* the compiler is the project's own clang, the same one that compiles the C++ + half, invoked `-x cuda` with `-D__HIP_PLATFORM_NVIDIA__`; +* `xim:hip-nvidia` contains no binaries, because on this platform there are + none to contain; +* the payloads are CUDA's, and `[dependencies.compat] cuda-runtime` is the same + one hop example 09 needs — mcpp's private loader does not consult `/usr/lib`, + so the statically linked CUDA runtime could not otherwise `dlopen` the driver. + +One of those payloads, `xim:cuda-profiler-api`, is here because CI found it +missing. `nvidia_hip_runtime_api.h` includes `` on its +second line and CUDA ships that header in a separate component, so a machine +with a host CUDA installation supplies it from `/usr/include` and the build +works while depending on something it never declared. + +`hipcc` is deliberately not used. It is a driver that reads `HIP_PLATFORM`, +picks nvcc or amdclang and forwards; every decision it makes is one the rule +has already made from the declaration, and it would make them again from the +environment. + +## The device is spelled once + +```toml +[build] +accel = "hip, cuda12.9+{sm_89}" +``` + +Two chunks. The first names the programming model, the second names the device +— and the second is character for character what example 09 writes. A device +has one spelling in this ecosystem however many programming models reach it, so +`sm_89` does not acquire a second one because the file is called `.hip`. + +## The AMD platform + +Refused, by name, with the reason. It needs a ROCm runtime and device library +that this ecosystem does not publish yet, so compiling for it would produce an +object nothing on the machine can link or run. That refusal is the honest +answer; the alternative is a build that succeeds and an artifact that does not. + +## Running it + +``` +mcpp build # sm_89, through mcpp.rules.hip +mcpp run # 12 24 36 48, on the device +mcpp build --no-accel # the constrained glob is left out +mcpp run --no-accel # 12 24 36 48, from src/cpu/saxpy.cpp +``` diff --git a/examples/12-hip-kernel/app/build.mcpp b/examples/12-hip-kernel/app/build.mcpp new file mode 100644 index 00000000..81f835ad --- /dev/null +++ b/examples/12-hip-kernel/app/build.mcpp @@ -0,0 +1,9 @@ +import std; +import mcpp; +import mcpp.rules.hip; + +int main() { + mcpp::rules::hip::options opt; + opt.includes = { "include" }; + return mcpp::rules::hip::compile(opt) ? 0 : 1; +} diff --git a/examples/12-hip-kernel/app/include/saxpy/saxpy.h b/examples/12-hip-kernel/app/include/saxpy/saxpy.h new file mode 100644 index 00000000..8d82d972 --- /dev/null +++ b/examples/12-hip-kernel/app/include/saxpy/saxpy.h @@ -0,0 +1,21 @@ +// 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); + +#ifdef __cplusplus +} +#endif +#endif diff --git a/examples/12-hip-kernel/app/mcpp.toml b/examples/12-hip-kernel/app/mcpp.toml new file mode 100644 index 00000000..c47885b9 --- /dev/null +++ b/examples/12-hip-kernel/app/mcpp.toml @@ -0,0 +1,70 @@ +[package] +name = "hip-saxpy" +namespace = "example" +version = "0.1.0" +description = "A HIP kernel behind a seam module, reaching an NVIDIA device, with a CPU fallback" +accelerators = ["hip"] + +[language] +standard = "c++23" +modules = true +import_std = true + +# The same clang that compiles the C++ half compiles the device unit. HIP on +# the NVIDIA platform is a header layer over the CUDA runtime, so there is no +# hipcc, no ROCm, and no second host compiler to satisfy a bound for. +[toolchain] +default = "llvm@22.1.8" + +[dependencies.mcpp] +plugins = { version = "0.2.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 +# private loader does not consult /usr/lib, and without a directory on the +# artifact's runtime search path the statically linked CUDA runtime cannot +# dlopen the driver. +[dependencies.compat] +cuda-runtime = "2026.09.05" + +# The payloads. `hip-nvidia` is headers only -- on this platform that is all +# HIP is -- and the four CUDA entries are the back end it compiles through. +# cuRAND and CCCL are on the list for the reason the CUDA consumer records: +# clang's CUDA wrapper includes curand_mtgp32_kernel.h for every device unit +# and that header includes , so a unit that calls neither still +# needs both. +[xlings.workspace] +"xim:hip-nvidia" = "7.2.4" +"xim:cuda-nvcc" = "12.9.86" +"xim:cuda-cudart" = "12.9.79" +"xim:libcurand" = "10.3.10.19" +"xim:cuda-cccl" = "12.9.27" +# `nvidia_hip_runtime_api.h` includes at its second line, +# and CUDA ships that header in its own component. A machine with a host CUDA +# installation finds it in /usr/include without saying so, which is how this +# entry came to be missing from a build that worked. +"xim:cuda-profiler-api" = "12.9.79" +"xim:libcuda-host-link" = { linux = "0.0.1" } + +[build] +# Two chunks: the programming model, and the device. A device is spelled once +# in this ecosystem however many models reach it, so `sm_89` here is the same +# `sm_89` the CUDA consumer writes. +accel = "hip, cuda12.9+{sm_89}" +sources = [ + "src/*.cppm", + "src/*.cpp", + { glob = "src/kernels/**/*.hip", accel = "hip, cuda12.9+{sm_89}" }, +] +include_dirs = ["include"] + +[target.'cfg(accelerator = "hip")'.build] +ldflags = ["-lcudart_static", "-lrt", "-lpthread", "-ldl"] + +# The CPU-only variant: the same seam, a host implementation behind it. +[target.'cfg(not(accelerator = "hip"))'.build] +sources = ["src/cpu/*.cpp"] + +[targets.hip-saxpy] +kind = "bin" +main = "src/main.cpp" diff --git a/examples/12-hip-kernel/app/src/app.cppm b/examples/12-hip-kernel/app/src/app.cppm new file mode 100644 index 00000000..a5e1a3ff --- /dev/null +++ b/examples/12-hip-kernel/app/src/app.cppm @@ -0,0 +1,27 @@ +// The seam. +// +// Its reason for existing is not that a device compiler rejects modules. It is +// that this is the one place a backend can be exchanged: the island underneath +// 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; + +export namespace app { + +// The device interface is raw pointers and a count because it has to be. The +// seam is where that becomes a C++ interface again. +std::optional> +saxpy(float a, std::span x, std::span y) { + if (x.size() != y.size()) return std::nullopt; + std::vector out(x.size()); + if (saxpy_device(a, x.data(), y.data(), out.data(), + static_cast(x.size())) != 0) + return std::nullopt; + return out; +} + +} // namespace app diff --git a/examples/12-hip-kernel/app/src/cpu/saxpy.cpp b/examples/12-hip-kernel/app/src/cpu/saxpy.cpp new file mode 100644 index 00000000..06a6ef0f --- /dev/null +++ b/examples/12-hip-kernel/app/src/cpu/saxpy.cpp @@ -0,0 +1,11 @@ +// The CPU implementation behind the same seam. Compiled only when the build +// asks for no accelerator (`mcpp build --no-accel`), through the +// `cfg(not(accelerator = "hip"))` section of the manifest; the device island +// and this file define the same symbol and are never in one link. +#include "saxpy/saxpy.h" + +extern "C" int saxpy_device(float a, const float* x, const float* y, + float* out, unsigned n) { + for (unsigned i = 0; i < n; ++i) out[i] = a * x[i] + y[i]; + return 0; +} diff --git a/examples/12-hip-kernel/app/src/kernels/saxpy.hip b/examples/12-hip-kernel/app/src/kernels/saxpy.hip new file mode 100644 index 00000000..8085df6a --- /dev/null +++ b/examples/12-hip-kernel/app/src/kernels/saxpy.hip @@ -0,0 +1,63 @@ +// The island, written against the HIP API. Nothing here is visible to the +// module graph: a device translation unit is never scanned and never produces +// a BMI. +// +// On the NVIDIA platform every name below is an inline wrapper over the CUDA +// one -- `hipMalloc` resolves to `cudaMalloc` through the header -- so this +// file and its CUDA sibling in tests/cuda-consumer compile to the same device +// code through different spellings. That is the point of the fixture: the +// engine path, the rule and the seam are the same, and only the API differs. +// +// It uses no C++ standard library. That is a deliberate property rather than +// an accident of a small example: an island that pulls in libstdc++ links a +// 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 + +namespace { + +__global__ void saxpy_kernel(float a, const float* x, const float* y, + float* out, unsigned n) { + unsigned i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) out[i] = a * x[i] + y[i]; +} + +} // namespace + +extern "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; + + if (hipError_t e = hipMalloc(&dx, bytes); e != hipSuccess) { + std::fprintf(stderr, "hipMalloc: %s\n", hipGetErrorString(e)); + goto done; + } + if (hipMalloc(&dy, bytes) != hipSuccess) goto done; + if (hipMalloc(&dout, bytes) != hipSuccess) goto done; + if (hipMemcpy(dx, x, bytes, hipMemcpyHostToDevice) != hipSuccess) goto done; + if (hipMemcpy(dy, y, bytes, hipMemcpyHostToDevice) != hipSuccess) goto done; + + hipLaunchKernelGGL(saxpy_kernel, dim3((n + 255) / 256), dim3(256), 0, 0, + a, dx, dy, dout, n); + // The launch is asynchronous, so its own return value reports only whether + // the launch was accepted. A kernel compiled for an architecture this + // device does not have fails HERE, with `no kernel image is available for + // execution on the device` — which is the runtime failure the accelerator + // dimension of an artifact's identity exists to turn into a build-time one. + if (hipError_t e = hipGetLastError(); e != hipSuccess) { + std::fprintf(stderr, "launch: %s\n", hipGetErrorString(e)); + goto done; + } + if (hipDeviceSynchronize() != hipSuccess) goto done; + if (hipMemcpy(out, dout, bytes, hipMemcpyDeviceToHost) != hipSuccess) goto done; + rc = 0; + +done: + hipFree(dx); hipFree(dy); hipFree(dout); + return rc; +} diff --git a/examples/12-hip-kernel/app/src/main.cpp b/examples/12-hip-kernel/app/src/main.cpp new file mode 100644 index 00000000..6efc1938 --- /dev/null +++ b/examples/12-hip-kernel/app/src/main.cpp @@ -0,0 +1,12 @@ +import std; +import app.saxpy; + +int main() { + const std::vector x{1, 2, 3, 4}, y{10, 20, 30, 40}; + auto out = app::saxpy(2.0f, x, y); + if (!out) { std::println("device unavailable"); return 1; } + for (auto v : *out) std::print("{} ", v); + std::println(""); + const std::vector want{12, 24, 36, 48}; + return *out == want ? 0 : 1; +}