// Compile CUDA device translation units and hand the objects to the link.

//

// WHY A RULE PACKAGE RATHER THAN THE ENGINE

//

// Everything below is knowledge about one vendor's tools: where the toolkit's

// pieces live, how an architecture is spelled, which host compilers nvcc

// tolerates, how a driver states its version. None of it is knowledge about

// the build graph. The engine owns the graph, the artifact's identity and the

// architecture set (`[build] accel`); the spelling of the command that

// produces an object, and every probe of the machine, is this file's business.

// `tests/unit/test_core_vendor_probes.cpp` in mcpp holds that line from the

// other side: the engine names no vendor tool.

//

// TWO ROUTES, ONE PRIMARY

//

// `clang -x cuda` is the primary route: the resolved toolchain's own clang

// compiles the device unit, so there is no second host compiler and no

// host-compiler bound to satisfy. nvcc is the alternate, taken when the

// project's toolchain is GCC or when asked for: it drives a host compiler

// (`-ccbin`) and refuses one newer than the bound its `crt/host_config.h`

// states, which this rule reads and reports.

//

// WHAT THE RULE TELLS THE ENGINE

//

// Objects, through `mcpp::action` (role "object"); library directories,

// through `mcpp::link_search`; and three claims about the machine that the

// engine compares or relays before the first compile:

// - the driver's version, read through the driver's own library and stated

// with `mcpp::fact`, together with the floor the runtime needs

// (`mcpp::floor`); an unmet floor refuses the build with both values;

// - whether nvcc can reach its own back-end stages (`--dryrun`), as an

// advisory naming the first stage that does not resolve;

// - whether the embedded PTX can be JIT-compiled by this driver, as an

// advisory, because the SASS for the named architectures still runs.

module;

#include <cstdlib>

#include <cstdio>

#if !defined(_WIN32)

#include <dlfcn.h>

#endif

export module mcpp.rules.cuda;

import std;

import mcpp;

import mcpp.plugins.tool;

// WHY NOTHING HERE USES `std::println`, AND WHY THAT IS NOT A STYLE CHOICE.

//

// `std::print` and `std::println` are not header-only. Both of their overloads

// reach into the libc++ DYLIB -- `__is_posix_terminal(FILE*)` for the stdout

// form and `__get_ostream_file(ostream&)` for the stream form -- and those

// symbols were added to that library in a version macOS 14 does not ship. A

// build program's link resolves `-lc++` to the system copy there, so a rule

// that printed with `std::println` compiled and then failed to link:

//

// ld64.lld: error: undefined symbol: std::__1::__is_posix_terminal(__sFILE*)

//

// naming neither the call that needed it nor the reason. Measured on

// macos-14; macos-15 has the symbol, which is why nothing saw this until a

// rule was first compiled on the older of the two supported releases.

//

// `std::format` is header-only and has no such dependency, so every message in

// this file is formatted and then streamed.

export namespace mcpp::rules::cuda {

// A PROGRAM'S NAME CARRIES ITS HOST'S SUFFIX, AND THE PAYLOAD LAYOUT DOES TOO.

//

// The NVIDIA redistributables are published for Windows as well as Linux and

// carry the same tree under different names: `bin/nvcc.exe` rather than

// `bin/nvcc`, and import libraries under `lib/x64` rather than shared objects

// under `lib` or `lib64`. Every one of the paths this rule builds is derived

// from those two facts, so both are stated once here rather than at each site

// -- a `std::filesystem::exists` against the wrong spelling does not fail, it

// answers false, and this rule reads several such answers as decisions.

#if defined(_WIN32)

inline constexpr bool kWindows = true;

inline constexpr const char* kExe = ".exe";

#else

inline constexpr bool kWindows = false;

inline constexpr const char* kExe = "";

#endif

enum class route { automatic, clang, nvcc };

struct options {

route which = route::automatic;

// THE TOOLKIT (0.19.0, mcpp#755): a root holding the whole of it

// (`mcpp::plugins::tool::root("/usr/local/cuda")`), `nvcc` itself, or the

// default -- the payloads this rule declares, or overrides of them. A

// toolkit named here is not downloaded: the payloads are declared on

// request.

mcpp::plugins::tool::choice toolkit;

// Header search paths for the island. Relative entries resolve against the

// package root; an ABSOLUTE entry is passed through unchanged.

//

// THE ABSOLUTE FORM IS FOR A DEPENDENCY'S HEADERS. A device compiler is

// a separate driver and inherits nothing from the C++ side's include

// configuration, so a package whose device code includes a dependency's

// header -- ggml's CUDA backend includes `cublas_v2.h` -- has to name that

// dependency's directory here, and it knows it only as the absolute path

// `mcpp::dep_dir` answered with.

std::vector<std::string> includes;

// FLAGS FOR THE ISLAND'S COMPILER, PASSED THROUGH UNEXAMINED.

//

// A device compiler is a separate driver with its own command line, and

// `mcpp::cflag`/`mcpp::cxxflag` reach mcpp's compiler rather than this one.

// The case this exists for is `mcpp.tools.island`, whose

// `force_include_flags` makes the island read its generated boundary header

// before its first line -- so the island names no generated file and the

// project has no header of its own. Project-wide flags cannot do that job:

// forcing a header into every C++ translation unit puts declarations ahead

// of `export module`, which no module interface unit accepts.

std::vector<std::string> flags;

std::string out_dir = std::string(mcpp::out_dir());

};

// ─── What the engine said ──────────────────────────────────────────────────

// The `cuda` chunk of `mcpp::accel()`, in this rule's own reading: the engine

// carries the string and compares it as a shape; what `sm_89` means is ours.

struct target {

std::string version; // "12.9"

std::vector<std::string> archs; // {"sm_89"}

std::string ptx; // "89" when a portable form is embedded

bool present = false;

};

// Split on one character. Written out rather than taken from <ranges>: GCC 16

// refuses the ranges split view instantiated inside an exported inline function

// when build.mcpp imports this module (`conflicting deduced return type for

// imported declaration ... view_interface::data()`), and clang does not.

inline std::vector<std::string_view> split(std::string_view s, char sep) {

std::vector<std::string_view> out;

for (std::size_t i = 0; i <= s.size();) {

auto j = s.find(sep, i);

out.push_back(s.substr(i, j == std::string_view::npos ? s.size() - i : j - i));

if (j == std::string_view::npos) break;

i = j + 1;

}

return out;

}

inline std::string_view trim(std::string_view s) {

while (!s.empty() && (s.front() == ' ' || s.front() == '\t')) s.remove_prefix(1);

while (!s.empty() && (s.back() == ' ' || s.back() == '\t')) s.remove_suffix(1);

return s;

}

inline target parse_target(std::string_view accel) {

target t;

for (std::size_t i = 0; i <= accel.size();) {

auto comma = accel.find(',', i);

auto open = accel.find('{', i), close = accel.find('}', i);

if (open != std::string_view::npos && close != std::string_view::npos

&& comma != std::string_view::npos && comma > open && comma < close)

comma = accel.find(',', close);

auto chunk = trim(comma == std::string_view::npos ? accel.substr(i)

: accel.substr(i, comma - i));

i = comma == std::string_view::npos ? accel.size() + 1 : comma + 1;

if (!chunk.starts_with("cuda")) continue;

t.present = true;

auto plus = chunk.find('+');

t.version = std::string(trim(chunk.substr(4, plus == std::string_view::npos

? chunk.size() - 4 : plus - 4)));

if (plus != std::string_view::npos) {

auto o = chunk.find('{', plus), c = chunk.find('}', plus);

if (o != std::string_view::npos && c != std::string_view::npos)

for (auto part : split(chunk.substr(o + 1, c - o - 1), ','))

if (auto a = trim(part); !a.empty()) t.archs.emplace_back(a);

auto tail = chunk.substr(c == std::string_view::npos ? chunk.size() : c + 1);

for (auto key : {"ptx>=", "floor>="})

if (auto p = tail.find(key); p != std::string_view::npos)

t.ptx = std::string(trim(tail.substr(p + std::string_view(key).size())));

}

}

return t;

}

// ─── This rule's share of the device sources ───────────────────────────────

//

// `mcpp::device_sources()` is the package's WHOLE device set, not this rule's

// share of it. A project with two backends puts a `.cu` and a `.comp` in one

// list, and every rule in that build program reads the same variable. Taking

// all of it works for exactly as long as a build has one rule in it, and then

// fails on the second -- not by dropping anything, but by handing a compiler a

// file it does not accept, with a message about that file's contents rather

// than about the rule that should have had it.

//

// So each rule takes the extensions it CLAIMS and leaves the rest to whoever

// claims those. A file no rule claims is not silently dropped either: the

// engine refuses a device source that reached no action, which is the one

// place that can see every rule's share at once.

constexpr std::string_view kClaimed[] = { ".cu" };

inline bool claims_extension(std::string_view path) {

const auto slash = path.find_last_of("/\\");

const auto name = slash == std::string_view::npos ? path : path.substr(slash + 1);

const auto dot = name.rfind('.');

if (dot == std::string_view::npos) return false;

const auto ext = name.substr(dot);

for (auto e : kClaimed) if (e == ext) return true;

return false;

}

inline std::vector<std::string> device_sources() {

std::vector<std::string> out;

for (auto part : split(std::string_view(mcpp::device_sources()), '\n'))

if (auto s = trim(part); !s.empty() && claims_extension(s)) out.emplace_back(s);

return out;

}

// ─── Locating the toolkit ──────────────────────────────────────────────────

// The toolkit this project declared under `[xlings.workspace]`, by component.

// The 13.x line splits the compiler across `cuda-nvcc`, `cuda-crt` and

// `libnvvm`; the 12.x line keeps them in `cuda-nvcc`. Either way the project

// names the compiler and this rule finds the pieces.

struct toolkit {

std::string nvcc_root, cudart_root, crt_root, driver_dir;

// cuRAND's headers. Not a runtime dependency of a kernel that never calls

// cuRAND: clang's own CUDA wrapper (`__clang_cuda_runtime_wrapper.h`)

// includes `curand_mtgp32_kernel.h` unconditionally, so the clang route

// cannot compile any device unit without them. Measured 2026-09-05: on a

// developer machine the header was found in the HOST's /usr/include and

// the leak went unnoticed until a runner with no host CUDA refused it.

// The same header then includes <nv/target> from CCCL (libcu++), which the

// 12.x toolkits ship as the separate `cuda-cccl` package, and which the

// host's /usr/include had supplied in the same way.

std::string curand_root, cccl_root;

std::string nvcc() const { return nvcc_root + "/bin/nvcc" + kExe; }

std::string host_config() const {

for (auto const* r : { &crt_root, &nvcc_root, &cudart_root }) {

if (r->empty()) continue;

auto p = *r + "/include/crt/host_config.h";

if (std::filesystem::exists(p)) return p;

}

return {};

}

std::vector<std::string> include_dirs() const {

std::vector<std::string> out;

for (auto const* r : { &cudart_root, &crt_root, &nvcc_root, &cccl_root, &curand_root })

if (!r->empty() && std::filesystem::is_directory(*r + "/include"))

out.push_back(*r + "/include");

// The 13.x CCCL payload nests its tree one directory down.

if (!cccl_root.empty() && std::filesystem::is_directory(cccl_root + "/include/cccl"))

out.push_back(cccl_root + "/include/cccl");

return out;

}

bool has_cccl() const {

return !cccl_root.empty()

&& (std::filesystem::exists(cccl_root + "/include/nv/target")

|| std::filesystem::exists(cccl_root + "/include/cccl/nv/target"));

}

std::vector<std::string> lib_dirs() const {

std::vector<std::string> out;

// `/lib/x64` is the Windows layout; it is searched on every host

// because a directory that does not exist contributes nothing, and a

// list keyed on the host would be a second place to keep in step.

for (auto const* r : { &cudart_root, &nvcc_root })

for (auto const* sub : { "/lib", "/lib64", "/lib/x64" })

if (!r->empty() && std::filesystem::is_directory(*r + sub))

out.push_back(*r + sub);

return out;

}

};

inline std::string xpkg(const char* name) {

if (const char* d = mcpp::xpkg_dir("xim", name); d && *d) return d;

return {};

}

// The components this rule declares, in the order they are asked for.

inline constexpr const char* kComponents[] = {

"cuda-nvcc", "cuda-cudart", "cuda-crt", "libcurand", "cuda-cccl", "libcuda-host-link",

};

inline const mcpp::plugins::tool::spec& nvcc_spec() {

static const mcpp::plugins::tool::spec s{

.who = "mcpp.rules.cuda", .package = "cuda-nvcc", .programs = {"nvcc"},

.option = "options::toolkit" };

return s;

}

// `pending` is set when the toolkit was asked for: the engine installs every

// component in one batch and runs the program again.

inline std::optional<toolkit> find_toolkit(const options& opt, bool& pending) {

pending = false;

if (!opt.toolkit.is_default()) {

// ONE ROOT FOR EVERY COMPONENT: an NVIDIA installation keeps nvcc,

// the runtime, CCCL and cuRAND in one tree. The driver library is the

// machine's, unless the payload that links against it is installed.

auto f = mcpp::plugins::tool::resolve(nvcc_spec(), opt.toolkit);

if (!f) {

std::cerr << mcpp::plugins::tool::describe_missing(nvcc_spec(), f) << '\n';

return std::nullopt;

}

toolkit t;

t.nvcc_root = t.cudart_root = t.crt_root = t.curand_root = t.cccl_root = f.root;

t.driver_dir = xpkg("libcuda-host-link");

return t;

}

for (auto const* c : kComponents)

if (std::string_view(mcpp::xpkg_source("xim", c)) == "pending") {

(void)mcpp::xpkg_request("xim", c);

pending = true;

}

if (pending) return std::nullopt;

(void)mcpp::plugins::tool::resolve(nvcc_spec()); // the decision record

toolkit t;

t.nvcc_root = xpkg("cuda-nvcc");

t.cudart_root = xpkg("cuda-cudart");

t.crt_root = xpkg("cuda-crt");

t.curand_root = xpkg("libcurand");

t.cccl_root = xpkg("cuda-cccl");

t.driver_dir = xpkg("libcuda-host-link");

if (t.nvcc_root.empty() || t.cudart_root.empty()) {

std::cerr << std::format("mcpp.rules.cuda: the toolkit is not installed.\n"

" This rule DECLARES it, so a project normally writes nothing. Three "

"things stop\n"

" that from reaching the build, in the order worth checking:\n"

" - mcpp older than 2026.9.6.6, which cannot answer a payload a "

"dependency declared;\n"

" - `features = [\"rules-cuda\"]` missing from the "

"[build-dependencies] edge;\n"

" - the build names no CUDA accelerator (`--accel \"cuda12.9+{{sm_89}}\"` "

"or [build] accel).\n"

" To pin a different line, name it in your own project and it wins:\n"

" [target.'cfg(accelerator = \"cuda\")'.xlings.workspace]\n"

" \"xim:cuda-nvcc\" = \"12.9.86\"\n"

" \"xim:cuda-cudart\" = \"12.9.79\"\n"

" (found nvcc: '{}', cudart: '{}')", t.nvcc_root, t.cudart_root) << '\n';

return std::nullopt;

}

return t;

}

// ─── Probes: what the machine has, what the toolkit needs ──────────────────

// The driver's version through the driver's own library, reached through the

// sentinel package rather than /usr/lib. "" when there is no driver here,

// which is a fact about the machine and not a failure of the build.

inline std::string driver_version(const toolkit& t) {

#if defined(_WIN32)

return {};

#else

if (t.driver_dir.empty()) return {};

const auto lib = t.driver_dir + "/lib/libcuda.so.1";

if (!std::filesystem::exists(lib)) return {};

// What would change the answer is the library the answer was read from.

mcpp::rerun_if_changed(lib.c_str());

void* h = ::dlopen(lib.c_str(), RTLD_LAZY | RTLD_LOCAL);

if (!h) {

const char* why = ::dlerror();

mcpp::warning(std::format("could not open the driver library {}: {}", lib,

why ? why : "(no reason given)").c_str());

return {};

}

using fn = int (*)(int*);

auto get = reinterpret_cast<fn>(::dlsym(h, "cuDriverGetVersion"));

int v = 0;

std::string out;

if (get && get(&v) == 0 && v > 0) out = std::format("{}.{}", v / 1000, (v % 1000) / 10);

::dlclose(h);

return out;

#endif

}

inline int major_of(std::string_view v) {

int m = 0;

for (char c : v) { if (!std::isdigit(static_cast<unsigned char>(c))) break; m = m * 10 + (c - '0'); }

return m;

}

inline bool version_at_least(std::string_view have, std::string_view want) {

auto parse = [](std::string_view s) {

std::vector<int> out; int acc = 0; bool digits = false;

for (char c : s) {

if (c == '.') { out.push_back(acc); acc = 0; digits = false; continue; }

if (!std::isdigit(static_cast<unsigned char>(c))) break;

acc = acc * 10 + (c - '0'); digits = true;

}

if (digits) out.push_back(acc);

return out;

};

auto h = parse(have), w = parse(want);

for (std::size_t i = 0; i < std::max(h.size(), w.size()); ++i) {

int a = i < h.size() ? h[i] : 0, b = i < w.size() ? w[i] : 0;

if (a != b) return a > b;

}

return true;

}

// State the driver relation. The engine compares the floor against the fact

// and refuses with both values; this rule only knows which numbers matter.

//

// The floor is the toolkit's major: a 12.x runtime runs on any 12.x driver

// (minor-version compatibility), and fails at the first allocation on an 11.x

// one. The embedded PTX is a separate, softer question: PTX emitted by toolkit

// 12.9 is JIT-compiled only by a driver at or above 12.9, but the SASS for the

// named architectures still runs, so a driver below the toolkit costs reach on

// newer hardware rather than correctness here -- reported, not enforced.

inline void state_driver_relation(const toolkit& t, const target& tg) {

const auto driver = driver_version(t);

if (!driver.empty()) mcpp::fact("cuda.driver", driver.c_str());

else mcpp::warning("no driver library reachable through xim:libcuda-host-link; "

"the build proceeds and the artifact will find no device at run time");

const int major = major_of(tg.version);

if (major > 0) mcpp::floor(std::format("cuda.driver >= {}.0", major).c_str());

if (!driver.empty() && !tg.ptx.empty() && !version_at_least(driver, tg.version))

mcpp::warning(std::format(

"the PTX embedded for compute_{} was emitted by toolkit {} and this driver "

"serves {}; hardware newer than {{{}}} will not be able to JIT it. The named "

"architectures run. Build with a toolkit at or below the driver, or add the "

"newer hardware's SASS to [build] accel.",

tg.ptx, tg.version, driver, [&] {

std::string s; for (auto& a : tg.archs) { if (!s.empty()) s += ','; s += a; }

return s; }()).c_str());

}

// The greatest gcc major and the greatest clang major the toolkit accepts.

// Zero means the header said nothing, which is not a refusal.

struct bounds { int gcc = 0, clang = 0; };

inline bounds read_bounds(std::string_view headerPath) {

bounds b;

if (headerPath.empty()) return b;

std::ifstream in{std::string(headerPath)};

std::string text{std::istreambuf_iterator<char>(in), std::istreambuf_iterator<char>()};

auto number_after = [&](std::size_t pos) {

int v = 0, n = 0;

while (pos < text.size() && !std::isdigit(static_cast<unsigned char>(text[pos]))) {

if (text[pos] == '\n') return 0;

++pos;

}

while (pos < text.size() && std::isdigit(static_cast<unsigned char>(text[pos]))) {

v = v * 10 + (text[pos] - '0'); ++pos; ++n;

}

return n ? v : 0;

};

if (auto p = text.find("__GNUC__ > "); p != std::string::npos) b.gcc = number_after(p + 10);

if (auto p = text.find("clang version must be less than "); p != std::string::npos)

if (int excl = number_after(p + 31); excl > 0) b.clang = excl - 1;

return b;

}

// Does the C library this build compiles against declare the C23 functions

// `cospi`, `sinpi` and `rsqrt`?

//

// Measured 2026-09-05 against glibc 2.44. Toolkit 12.9's

// `crt/math_functions.h` declares those same names for the host WITHOUT

// `noexcept`; glibc declares them WITH it, and since C++17 that is part of the

// function type. nvcc's front end stops with six `exception specification is

// incompatible` errors that name a glibc header and a CUDA header and leave

// the reader to work out that neither is at fault alone. The 13.x line does

// not redeclare them and compiles cleanly against the same C library.

//

// Read, not probed. The answer is one substring of one header the sysroot

// already contains; a probe compile would spend a second nvcc invocation to

// learn the same thing, and would report it as a compile failure rather than

// as a pairing that cannot work.

inline bool libc_declares_c23_pi_math(std::string_view sysroot) {

if (sysroot.empty()) return false;

for (auto const* rel : { "/usr/include/bits/mathcalls.h", "/include/bits/mathcalls.h" }) {

std::ifstream in{std::string(sysroot) + rel};

if (!in) continue;

std::string text{std::istreambuf_iterator<char>(in), std::istreambuf_iterator<char>()};

return text.find("(cospi,") != std::string::npos

&& text.find("(rsqrt,") != std::string::npos;

}

return false;

}

// `<compiler> -dumpversion` → major. The host compiler is the toolchain mcpp

// resolved for this build, so its version is a fact of the build, not a guess.

inline int compiler_major(const std::string& cc) {

#if defined(_WIN32)

(void)cc; return 0;

#else

std::string cmd = cc + " -dumpversion 2>/dev/null";

if (FILE* p = ::popen(cmd.c_str(), "r")) {

char buf[64] = {};

std::string s;

if (std::fgets(buf, sizeof buf, p)) s = buf;

::pclose(p);

return major_of(s);

}

return 0;

#endif

}

// The first back-end stage nvcc names but cannot resolve, from its own plan.

// nvcc invokes cicc, cudafe++, ptxas and fatbinary by bare name on a PATH it

// states in the plan; a stage that does not resolve there fails the compile

// with `sh: 1: cicc: not found`, naming nothing that helps.

inline std::optional<std::string> unreachable_stage(const toolkit& t, const std::string& ccbin) {

#if defined(_WIN32)

(void)t; (void)ccbin; return std::nullopt;

#else

const auto probe = std::filesystem::temp_directory_path() / "mcpp-rules-cuda-dryrun.cu";

{ std::ofstream(probe) << "__global__ void k() {}\n"; }

std::string cmd = std::format("{} --dryrun -ccbin {} -c {} -o /dev/null 2>&1",

t.nvcc(), ccbin, probe.string());

std::string text;

if (FILE* p = ::popen(cmd.c_str(), "r")) {

char buf[4096];

while (std::fgets(buf, sizeof buf, p)) text += buf;

::pclose(p);

}

std::filesystem::remove(probe);

std::string path;

std::vector<std::string> stages;

for (auto l : split(text, '\n')) {

if (!l.starts_with("#$ ")) continue;

l.remove_prefix(3);

if (l.starts_with("PATH=")) { path = std::string(l.substr(5)); continue; }

for (auto const* stage : { "cicc", "cudafe++", "ptxas", "fatbinary", "nvlink" }) {

auto pos = l.find(stage);

if (pos == 0 || (pos != std::string_view::npos && (l[pos - 1] == ' ' || l[pos - 1] == '"')))

if (std::ranges::find(stages, stage) == stages.end()) stages.emplace_back(stage);

}

}

if (stages.empty()) return std::nullopt; // no plan, no finding

for (auto const& stage : stages) {

bool found = false;

for (auto dir : split(path, ':')) {

std::string d(dir);

if (!d.empty() && std::filesystem::exists(d + "/" + stage)) { found = true; break; }

}

if (!found) return stage;

}

return std::nullopt;

#endif

}

// ─── Planning ──────────────────────────────────────────────────────────────

struct edge {

std::string id, description;

std::vector<std::string> command, inputs, outputs;

// A depfile the COMMAND writes and ninja reads back. `inputs` is fixed

// when this program runs, before the compiler has seen the source, so a

// device unit that `#include`s a header had no edge to it: editing the

// header rebuilt nothing and the build stayed green over a stale object.

// The compiler already computes the answer while parsing.

//

// `-MMD` and not `-MD`, which is the choice mcpp makes for its own C and

// GAS units: user includes only. `-MD` was measured on BiSheng and pulled

// in fifty host headers under /usr/include -- correct, and useless, because

// it makes the object depend on absolute host paths that a shared build

// directory must not carry. A toolkit header is not missed by this: the

// toolkit's path carries its version, so a different toolkit is a different

// command line, which ninja already tracks.

std::string depfile;

};

// ON WINDOWS THE DEFAULT IS THE CLANG ROUTE WHATEVER THE PROJECT'S COMPILER IS.

//

// The two routes differ in who compiles the host half of a device unit. On the

// nvcc route that is a compiler nvcc drives through `-ccbin`, and on Windows

// the only one it accepts is MSVC's `cl.exe` -- whose directory is found

// through `vswhere` and the installed Visual Studio instance, neither of which

// mcpp resolves or states to a build program. The rule would have to search

// the host for it, which is the one thing a rule of this ecosystem does not do.

//

// The clang route has no such gap: it drives no second compiler, and clang

// locates the MSVC headers and libraries it needs for the host half itself,

// the same way it does for every ordinary C++ translation unit on this host.

inline route decide(route asked) {

if (asked != route::automatic) return asked;

#if defined(_WIN32)

return route::clang;

#else

return std::string_view(mcpp::compiler()) == "clang" ? route::clang : route::nvcc;

#endif

}

inline std::vector<edge> plan(std::span<const std::string> sources, options opt = {}) {

std::vector<edge> out;

const std::string root = mcpp::manifest_dir();

if (root.empty()) {

std::cerr << std::format("mcpp.rules.cuda: no mcpp build context -- this runs from build.mcpp") << '\n';

return out;

}

const auto tg = parse_target(mcpp::accel());

if (!tg.present || tg.archs.empty()) {

// C19: a device build that names no device is refused HERE, not at

// run time as `no kernel image is available for execution`.

std::cerr << std::format("mcpp.rules.cuda: [build] accel names no CUDA architecture (accel = \"{}\").\n"

" Write e.g. accel = \"cuda12.9+{{sm_89}} ptx>=89\" -- the set a build compiles\n"

" for is a decision, and the machine's own hardware is a poor default for it.",

mcpp::accel()) << '\n';

return out;

}

bool tkPending = false;

auto tk = find_toolkit(opt, tkPending);

if (!tk) return out;

state_driver_relation(*tk, tg);

const route r = decide(opt.which);

const std::string tcdir = mcpp::toolchain_dir();

std::string driver_cc; // the compiler that runs the device unit

std::vector<std::string> front; // the command up to the input file

if (r == route::clang) {

driver_cc = tcdir + "/bin/clang++" + kExe;

if (!std::filesystem::exists(driver_cc)) {

std::cerr << std::format("mcpp.rules.cuda: the clang route needs the toolchain's clang++ at {}.\n"

" Name an LLVM toolchain for this project:\n"

" [toolchain]\n"

" default = \"llvm@22.1.8\"{}",

driver_cc,

kWindows ? "\n On Windows this route is the only one the rule takes: nvcc "

"drives a host\n compiler named by -ccbin, and locating MSVC's "

"cl.exe is not something\n this rule does."

: "") << '\n';

return out;

}

// Refused here rather than at clang's include error: the header it

// would fail on belongs to a payload the project has to name, and the

// diagnostic names it. A host copy is never searched for -- that is

// how the leak above survived every local build.

const bool curand_ok = !tk->curand_root.empty()

&& std::filesystem::exists(tk->curand_root + "/include/curand_mtgp32_kernel.h");

if (!curand_ok || !tk->has_cccl()) {

std::cerr << std::format("mcpp.rules.cuda: the clang route needs cuRAND's headers, which clang's CUDA "

"wrapper includes unconditionally, and CCCL's, which they include in turn.\n"

" This rule declares both; see the note above for why they may not have "

"arrived.\n"

" To pin a different line, name it in your own project and it wins:\n"

" \"xim:cuda-cccl\" = \"12.9.27\" (the 12.9 line; 13.x pairs with 13.x)\n"

" \"xim:libcurand\" = \"10.3.10.19\" (the 12.9 line; 10.4.x pairs with 13.x)\n"

" (found cccl: '{}', curand: '{}')", tk->cccl_root, tk->curand_root) << '\n';

return out;

}

front = { driver_cc, "-x", "cuda", "-std=c++17", "-O2",

"--cuda-path=" + tk->nvcc_root, "-Wno-unknown-cuda-version",

// NVIDIA'S HEADER REFUSES libc++, AND THE REFUSAL IS

// ABOUT nvcc RATHER THAN ABOUT THIS COMPILER.

//

// crt/host_defines.h:67: error: "libc++ is not supported

// on x86 system"

//

// The guard is `#if defined(__CUDACC__) && … &&

// defined(_LIBCPP_VERSION)`, and clang defines `__CUDACC__`

// when it compiles CUDA itself — so a device unit that

// includes <cuda_runtime.h> stops here on any LLVM

// toolchain, which is the toolchain this route exists for.

// Measured on ggml's CUDA backend; the CUDA example's own

// kernel never showed it because a bare kernel includes no

// toolkit header at all.

//

// The escape hatch is upstream's own, and it is passed only

// on this route: nvcc's host pass really does break against

// libc++, and nothing here weakens that.

"-D_ALLOW_UNSUPPORTED_LIBCPP" };

// Position-independent code is the default on Windows and naming it

// is an unused-argument warning on every device unit.

if constexpr (!kWindows) front.insert(front.begin() + 5, "-fPIC");

for (auto const& inc : tk->include_dirs()) front.push_back("-I" + inc);

for (auto const& a : tg.archs) front.push_back("--cuda-gpu-arch=" + a);

// clang checks ptxas and fatbinary itself; say so before it does.

for (auto const* tool : { "ptxas", "fatbinary" })

if (!std::filesystem::exists(tk->nvcc_root + "/bin/" + tool + kExe))

mcpp::warning(std::format("the toolkit payload has no {}; clang invokes it "

"after generating PTX", tool).c_str());

std::cout << std::format("mcpp.rules.cuda: clang route -- {} (toolkit {})", driver_cc, tk->nvcc_root) << '\n';

} else {

// ASKED FOR EXPLICITLY, BECAUSE `decide` NEVER CHOOSES IT HERE.

// Everything below reads a GCC bound out of the toolkit's header and

// hands nvcc a `g++` to drive. On Windows there is no such compiler in

// the toolchain and the bound in `host_config.h` is stated in

// `_MSC_VER`, so the branch would build a command line out of two

// answers that mean nothing and nvcc would report the third.

if constexpr (kWindows) {

std::cerr << std::format("mcpp.rules.cuda: the nvcc route was asked for, and on Windows this rule does "

"not take it.\n"

" nvcc compiles the host half through a compiler named by -ccbin, which on "

"this host\n is MSVC's cl.exe; finding it means asking the machine about its "

"Visual Studio\n installation, which is the kind of host dependence this "

"ecosystem removes.\n"

" Leave the route unset: the clang route is the default here and needs no "

"second compiler.") << '\n';

return out;

}

// nvcc drives the toolchain's own compiler, and refuses one newer than

// the bound its header states. Read the bound; if exceeded, pass the

// escape hatch and say so -- an unexplained flag is worse than a note.

const bool clangHost = std::string_view(mcpp::compiler()) == "clang";

if (clangHost) {

// Measured: nvcc's own crt/host_defines.h stops the compile with

// `libc++ is not supported on x86 system`, and libc++ is what an

// LLVM toolchain's clang uses. The pairing that works is nvcc with

// a GCC toolchain; with an LLVM toolchain the clang route is the

// one to take, and it is the default.

std::cerr << std::format("mcpp.rules.cuda: the nvcc route needs a GCC host compiler; this project's "

"toolchain is LLVM, whose clang uses libc++ and nvcc refuses it. Use the clang "

"route (the default for an LLVM toolchain) or set [toolchain] to a gcc payload.") << '\n';

return out;

}

// The other pairing this route cannot have: an old toolkit and a C

// library new enough to have the C23 `pi` functions. Stated before the

// compile, because the compile's own report names two headers and no

// decision.

if (major_of(tg.version) < 13

&& libc_declares_c23_pi_math(mcpp::toolchain_sysroot())) {

std::cerr << std::format("mcpp.rules.cuda: toolkit {} redeclares the C23 functions cospi, sinpi and "

"rsqrt for the host without `noexcept`, and the C library this build compiles "

"against declares them with it; nvcc's front end refuses the pair.\n"

" Name a 13.x toolkit, whose headers leave them to the C library:\n"

" [xlings.workspace]\n"

" \"xim:cuda-nvcc\" = \"13.3.33\"\n"

" \"xim:cuda-crt\" = \"13.3.33\"\n"

" \"xim:cuda-cudart\" = \"13.3.29\"\n"

" or take the clang route, which does not include that header at all.",

tg.version) << '\n';

return out;

}

// The host compiler nvcc drives, chosen within the bound the toolkit

// states. Measured: gcc 16 under nvcc 12.9 (bound gcc <= 14) fails inside

// nvcc's front end on GCC 16's own <type_traits> even with

// -allow-unsupported-compiler -- the escape hatch admits a compiler one

// step past the bound, not a standard library two majors newer. So the

// rule does not guess: the toolchain's g++ when it is within the bound,

// otherwise a gcc payload the project declared for this purpose, and

// otherwise a refusal that says which declaration to add.

const auto b = read_bounds(tk->host_config());

const std::string tcGcc = tcdir + "/bin/g++" + kExe;

const int tcMajor = compiler_major(tcGcc);

if (b.gcc == 0 || tcMajor <= b.gcc) {

driver_cc = tcGcc;

} else if (auto payload = xpkg("gcc"); !payload.empty()

&& compiler_major(payload + "/bin/g++" + kExe) <= b.gcc) {

driver_cc = payload + "/bin/g++" + kExe;

mcpp::warning(std::format(

"nvcc {} states gcc <= {} in {}; the toolchain's gcc {} exceeds it, so the "

"device unit is compiled with the declared xim:gcc payload ({}). The clang "

"route has no such bound.", tg.version, b.gcc, tk->host_config(), tcMajor,

driver_cc).c_str());

} else {

std::cerr << std::format("mcpp.rules.cuda: nvcc {} accepts gcc <= {} ({}), and this project's "

"toolchain is gcc {}.\n"

" Declare a gcc payload within the bound and the rule drives that one:\n"

" [xlings.workspace]\n"

" \"xim:gcc\" = \"13.3.0\"\n"

" or take the clang route with [toolchain] default = \"llvm@22.1.8\".",

tg.version, b.gcc, tk->host_config(), tcMajor) << '\n';

return out;

}

front = { tk->nvcc(), "-ccbin", driver_cc, "-std=c++17", "-O2",

"--compiler-options", "-fPIC" };

// The host compiler nvcc drives is not one mcpp resolved, so nothing

// has told it where the C library or the assembler are. Measured: with

// neither of these, NVIDIA's own crt/host_config.h stops at

// `features.h: No such file or directory`. Both are the flags mcpp

// passes to its own compiler for this target.

if (const char* sr = mcpp::toolchain_sysroot(); sr && *sr) {

front.push_back("--compiler-options");

front.push_back(std::string("--sysroot=") + sr);

}

if (const char* bu = mcpp::toolchain_binutils_dir(); bu && *bu) {

front.push_back("--compiler-options");

front.push_back(std::string("-B") + bu);

}

for (auto const& inc : tk->include_dirs()) front.push_back("-I" + inc);

for (auto const& a : tg.archs) {

std::string digits;

for (char c : a) if (std::isdigit(static_cast<unsigned char>(c))) digits += c;

front.push_back("-gencode");

front.push_back(std::format("arch=compute_{},code={}", digits, a));

}

if (!tg.ptx.empty()) {

front.push_back("-gencode");

front.push_back(std::format("arch=compute_{0},code=compute_{0}", tg.ptx));

}

if (auto missing = unreachable_stage(*tk, driver_cc))

mcpp::warning(std::format(

"nvcc cannot reach its own back-end: it invokes '{}' by name and that name "

"does not resolve on the search path it states. On the 13.x line install "

"xim:libnvvm beside xim:cuda-nvcc.", *missing).c_str());

std::cout << std::format("mcpp.rules.cuda: nvcc route -- {} with -ccbin {}", tk->nvcc(), driver_cc) << '\n';

}

// The link line gets its directories from here, not from the manifest: the

// rule resolved the payload, so the rule names where its libraries are.

for (auto const& d : tk->lib_dirs()) mcpp::link_search(d.c_str());

for (auto const& src : sources) {

const auto stem = std::filesystem::path(src).stem().string();

const auto obj = opt.out_dir + "/" + stem + ".cu.o";

edge e;

e.id = "cuda:" + stem;

e.description = (r == route::clang ? "clang -x cuda " : "nvcc ") + src;

e.command = front;

for (auto const& inc : opt.includes)

e.command.push_back("-I" + (std::filesystem::path(inc).is_absolute()

? inc : root + "/" + inc));

for (auto const& f : opt.flags) e.command.push_back(f);

// Where the compiler reports what it read. `<object>.d`, so two

// architectures of one source do not share a file and overwrite each

// other's answer -- the object name already carries the arch for

// exactly that reason.

e.depfile = obj + ".d";

e.command.insert(e.command.end(), { "-MMD", "-MF", e.depfile });

e.command.insert(e.command.end(), { "-c", root + "/" + src, "-o", obj });

e.inputs = { root + "/" + src };

e.outputs = { obj };

out.push_back(std::move(e));

}

return out;

}

inline bool submit(std::span<const edge> edges) {

for (auto const& e : edges) {

mcpp::action a;

a.id = e.id.c_str();

a.role = "object"; // the linkable artifact itself

a.description = e.description.c_str();

for (auto const& c : e.command) a.arg(c.c_str());

for (auto const& i : e.inputs) a.input(i.c_str());

for (auto const& o : e.outputs) a.output(o.c_str());

// Empty for an edge that declares none, which serialises identically to

// an action from before the field existed -- see mcpp::action::depfile.

if (!e.depfile.empty()) a.depfile = e.depfile.c_str();

a.submit();

}

return true;

}

// Everything from the manifest: the architectures from `[build] accel`, the

// sources from the constrained glob in `[build] sources`. A build that asks

// for no accelerator has no device sources and nothing to do here -- that is

// the CPU-only variant, and the seam's fallback carries it.

inline bool compile(options opt = {}) {

if (!*mcpp::accel()) return true;

// Several rules in one build program is the ordinary shape for a project

// with several backends, and each is called unconditionally -- the build

// program cannot know which backends this build named without parsing

// `accel` itself, which is what the rule already does. A rule whose

// backend this build does not name has nothing to do, and that is not a

// mistake and must not be reported as one.

if (!parse_target(mcpp::accel()).present) return true;

const auto sources = device_sources();

if (sources.empty()) {

mcpp::warning("[build] accel names cuda but no constrained glob matched a `.cu`; "

"nothing was compiled for it");

return true;

}

auto edges = plan(sources, std::move(opt));

if (edges.empty()) return false;

return submit(edges);

}

} // namespace mcpp::rules::cuda