route which = route::automatic;

// 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());

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;

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;

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;

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;

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;

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;

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;

}

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;

#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

int m = 0;

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

return m;

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;

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());

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;

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;

#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

#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

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;

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

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;

}

auto tk = find_toolkit();

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 = \"[email protected]\"{}",

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 = \"[email protected]\".",

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;

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;

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);