From 0d40731cbfdf61137f9b2968d76a4dfc5b4074fe Mon Sep 17 00:00:00 2001 From: Yael Date: Thu, 20 Aug 2026 17:20:22 +0000 Subject: [PATCH] fix(kernels): utils.h helpers must be host+device, not device-only `impl/kernels/utils.h` declared its whole first half -- swap, square, sqrt, pow, min, max, abs, sign, mod, typed_prod, prod, fillfrom, fill -- as `inline FF_CUDEV`, i.e. `__device__` only. Host code calls them: * `canUse32BitIndexMath`, in this same header, calls `typed_prod`. That is the edge behind every `FF_CANUSE32BITS` in `src/lib-cuda`. * Every `FF_CUHOST` launcher in `impl/cuda/{reg_field,reg_flow, distance_euclidean,distance_l1,distance_mesh}.h` calls `prod(size, n)` on its first line, to size the grid. * `impl/kernels/distance/mesh.h`'s `FF_CUHOST build_tree` calls `max`. nvcc rejects a host->device call only when the calling function is not itself a template. All of the above are templates, so nvcc emitted no diagnostic and cudafe++ wrote `{int volatile ___ = 1; ...; ::exit(___);}` into the HOST object in place of each callee's body. The result links cleanly, passes `-Wl,--no-undefined` and `ldd -r` (the damage is intra-TU), terminates the process with status 1 on the first call, and -- `exit` being `noreturn` -- loses every statement after that call at -O1 and above. Make the whole header `FF_CUHOSTDEV`. Nothing in it is device-specific; device codegen is unchanged and the host side gains one inline function per used instantiation. Without nvcc both macros expand to nothing, so the CPU layers preprocess byte-for-byte identically (verified over all ten `src/lib-cpu/*.cpp`). Two gates, because the fix alone leaves the trap set: * `tests/impl-cuda/compile_probe_hostdev.cu` calls each helper from a plain, non-template `__host__` function -- the shape nvcc does check -- so a re-qualified helper is a compile error again. It produces 22 errors against the pre-fix header and compiles in ~2 s. It also explicitly instantiates one launcher per affected header, which type-checks the launcher body in the device pass where the call edge is reported. * `tools/check-cuda-host-stubs.sh`, run in `build-cuda`, fails on any undefined `exit` in `build/obj/lib-cuda/*.o`. Nothing here calls `::exit`, so that symbol has exactly one source. `compile_probe_mesh.cu` documented this exact nvcc error as "a spurious artifact of the probe technique, not a defect". It was neither; that comment is corrected in place. Fixes #150. --- .github/workflows/ci.yml | 15 +++ CLAUDE.md | 15 +++ MIGRATION.md | 60 +++++++++++ include/fastfields/impl/kernels/utils.h | 84 ++++++++++------ tests/impl-cuda/compile_probe_hostdev.cu | 122 +++++++++++++++++++++++ tests/impl-cuda/compile_probe_mesh.cu | 37 +++++-- tools/check-cuda-host-stubs.sh | 68 +++++++++++++ 7 files changed, 364 insertions(+), 37 deletions(-) create mode 100644 tests/impl-cuda/compile_probe_hostdev.cu create mode 100755 tools/check-cuda-host-stubs.sh diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 81298d0..bc3a526 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -648,6 +648,21 @@ jobs: ${{ matrix.indexflags && format('INDEXFLAGS="{0}"', matrix.indexflags) || '' }} \ NVCC="/usr/bin/time -f 'FFMEM %M kB %e s %C' nvcc" \ 2>&1 | tee /tmp/build-cuda.log + - name: No host-pass exit(1) stubs in the CUDA objects + # fastfields-lib#150. nvcc diagnoses a host->device call only when the + # calling function is not a template; every host caller here is one, so + # instead of an error cudafe++ writes `::exit(1)` into the HOST object + # in place of the __device__-only callee's body, and -O1 then deletes + # everything after the call. Nothing else in this job can see that: the + # damage is intra-TU, so --no-undefined and `ldd -r` both pass, the + # objects keep their full size (front-end instantiation is unaffected), + # and the FFMEM budget below does not move either. + # + # An undefined `exit` in one of these objects has exactly one source -- + # that stub. Nothing in the tree calls `::exit`. + # + # Runs before the hub link so that a red here is unambiguous. + run: ./tools/check-cuda-host-stubs.sh build/obj/lib-cuda/*.o - name: Link the hub against the CUDA backend # The actual --no-undefined gate: libfastfields.so is what references # the FF_CUDA:: entry points, so this is where a module missing from diff --git a/CLAUDE.md b/CLAUDE.md index 843d61d..02543a1 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -184,6 +184,21 @@ pushpull's fully-static order×bound compile is nightly table lives above `MODULES` in `src/lib-cuda/Makefile`; read it before changing `-j`, `-O`, the bound/spline policy, or anything that makes a regulariser heavier. Do not recombine or "tidy" the split. +- **`FF_CUDEV` means "cannot run on the host", not "called from a kernel".** + nvcc diagnoses a `__host__` → `__device__` call **only when the calling + function is not itself a template**; every host caller in this tree is one + (the `FF_CUHOST` launchers in `impl/cuda/`, `canUse32BitIndexMath`, mesh.h's + `build_tree`). In that case it emits no diagnostic at all and `cudafe++` + replaces the callee's body *in the host object* with `::exit(1)` — which + links cleanly, passes `--no-undefined` and `ldd -r`, terminates the process + at the first call, and (because `exit` is `noreturn`) makes `-O1` delete + every statement after it. That shipped in `libfastfields-cuda.so` for months + with every job green; see fastfields-lib#150 and the qualifier table in + `MIGRATION.md`. **Everything in `impl/kernels/utils.h` is `FF_CUHOSTDEV` and + must stay that way**; two gates enforce it, the compile-time + `tests/impl-cuda/compile_probe_hostdev.cu` and the object-level + `tools/check-cuda-host-stubs.sh` in `build-cuda`. `FF_CUHOSTDEV` is the safe + default: identical device codegen, one extra inline host function. - **Macros in installed headers are `FF_`-prefixed.** Anything `#define`d under `include/` and not `#undef`'d before the end of that header is inherited by every downstream translation unit, so it must be namespaced by prefix. Macros diff --git a/MIGRATION.md b/MIGRATION.md index 9530a16..5e64062 100644 --- a/MIGRATION.md +++ b/MIGRATION.md @@ -231,6 +231,66 @@ per call, while **363 of the 391** `impl/cuda` upload sites use the enable, and its cost may partly offset the register-pressure win. Measured numbers and the exact benchmark that would settle it: #94 and #144. +## `__host__` / `__device__` qualifiers under nvcc (#150) + +**nvcc checks a host→device call only when the calling function is not itself +a template.** That single sentence is the whole hazard, and it is not +documented anywhere in the CUDA guide as a limitation. + +| caller | callee | nvcc 12.0 | +| --- | --- | --- | +| plain `__host__` function | `__device__` function | **error** | +| plain `__host__` function | `__device__` function *template* | **error** | +| `__host__` function *template* | `__device__` function | **accepted silently** | +| `__host__` function *template* | `__device__` function *template* | **accepted silently** | + +In the two bottom rows nvcc does not diagnose anything. `cudafe++` instead +emits, into the **host** object, this in place of the callee's real body (the +real one is kept next to it under `#if 0`): + +```c +{int volatile ___ = 1; (void)args; ::exit(___);} +``` + +So the caller compiles, links, passes `-Wl,--no-undefined` and `ldd -r` — the +damage is intra-TU, nothing becomes undefined — and **terminates the process +with status 1** the first time it runs. `exit` is `noreturn`, so from `-O1` +upwards the host compiler additionally deletes every statement after the call +as unreachable. `-O0` keeps those statements; it is not any less broken, the +process still exits. + +Every host caller in this tree is a template, which is why this went unseen: + +* `impl/kernels/utils.h`'s `canUse32BitIndexMath` → `typed_prod`, the edge + behind every `FF_CANUSE32BITS` in `src/lib-cuda`; +* every `FF_CUHOST` launcher in `impl/cuda/{reg_field,reg_flow, + distance_euclidean,distance_l1,distance_mesh}.h` → `prod(size, n)`, on its + first line, to size the grid; +* `impl/kernels/distance/mesh.h`'s `FF_CUHOST build_tree` → `max`. + +**The rule.** Anything in `impl/kernels/utils.h` is `FF_CUHOSTDEV`, without +exception — the header holds only backend-agnostic helpers over scalars and +raw arrays, and host code calls them. More generally, `FF_CUDEV` means "this +genuinely cannot run on the host" (a device intrinsic, `__shared__`, +`threadIdx`), not "this is called from a kernel". When in doubt, `FF_CUHOSTDEV` +costs nothing: device codegen is identical and the host copy is one inline +function. + +**The two gates**, added with the fix: + +1. `tests/impl-cuda/compile_probe_hostdev.cu` — calls each helper from a + plain, non-template `__host__` function, i.e. the shape in the top two rows + of the table, so a re-qualified helper is an nvcc **error** again. It also + explicitly instantiates one representative launcher per affected header, + which type-checks the whole launcher body (an explicit instantiation + instantiates it in the device pass, where the call edge *is* reported). + Compiles in ~2 s in `compile-probe-cuda`; against the pre-fix header it + produces 22 errors. +2. `tools/check-cuda-host-stubs.sh` — run in `build-cuda` over + `build/obj/lib-cuda/*.o`, fails on any undefined `exit`. Nothing in this + codebase calls `::exit`, so that symbol has exactly one source: the stub + above. This is the catch-all for edges the probe does not enumerate. + ## Porting pattern (per module) Use `distance.{h,cpp}` at each level as the template. diff --git a/include/fastfields/impl/kernels/utils.h b/include/fastfields/impl/kernels/utils.h index 37394ce..2c13bdd 100755 --- a/include/fastfields/impl/kernels/utils.h +++ b/include/fastfields/impl/kernels/utils.h @@ -15,6 +15,34 @@ FF_NAMESPACE_BEGIN(FF_NS) FF_NAMESPACE_BEGIN(FF_DEVICE) +// ============================================================================ +// EVERY function in this header is FF_CUHOSTDEV (`__host__ __device__`), and +// that is a contract, not an accident. Do not "tighten" one back to FF_CUDEV. +// +// Nothing here is device-specific -- these are small generic helpers over +// scalars and raw arrays -- and host code genuinely calls them: +// * `canUse32BitIndexMath` (below) calls `typed_prod`; +// * every FF_CUHOST launcher in impl/cuda/{reg_field,reg_flow,distance_*}.h +// sizes its grid with `prod(size, n)`; +// * `distance/mesh.h`'s FF_CUHOST `build_tree` calls `max`. +// +// When one of these was `__device__`-only (fastfields-lib#150), nvcc did NOT +// reject those calls. It diagnoses a host->device call only when the CALLER +// is a non-template function; every caller above is a template, so the check +// was skipped and cudafe++ emitted, into the HOST object, a body of +// +// {int volatile ___ = 1; (void)args; ::exit(___);} +// +// in place of the real one. The callers therefore compiled, linked, and +// passed `--no-undefined` and `ldd -r`, and terminated the process with +// status 1 at the first call -- at -O1 and above the host compiler also +// deleted everything after it, because `exit` is `noreturn`. +// +// tests/impl-cuda/compile_probe_hostdev.cu turns that silence back into a +// compile error: it calls each helper from a plain (non-template) host +// function, which is the shape nvcc does diagnose. +// ============================================================================ + // static check for floating types template struct is_floating_point { static constexpr bool value = false; }; @@ -29,14 +57,14 @@ struct is_floating_point { static constexpr bool value = true; }; template -inline FF_CUDEV +inline FF_CUHOSTDEV void swap(T& a, T& b) { T c(a); a=b; b=c; } template -inline FF_CUDEV +inline FF_CUHOSTDEV T square(T a) { return a*a; @@ -45,26 +73,26 @@ T square(T a) #ifdef __CUDACC__ template -inline FF_CUDEV +inline FF_CUHOSTDEV T sqrt(T a) {} template <> -inline FF_CUDEV +inline FF_CUHOSTDEV float sqrt(float a) { return ::sqrtf(a); } template <> -inline FF_CUDEV +inline FF_CUHOSTDEV double sqrt(double a) { return ::sqrt(a); } template <> -inline FF_CUDEV +inline FF_CUHOSTDEV half sqrt(half a) { // hsqrt is not visible at global scope in every CUDA/arch combination; @@ -75,7 +103,7 @@ half sqrt(half a) #else template -inline FF_CUDEV +inline FF_CUHOSTDEV T sqrt(T a) { return std::sqrt(a); @@ -85,7 +113,7 @@ T sqrt(T a) template -inline FF_CUDEV +inline FF_CUHOSTDEV T pow(T a) { T p = a; # pragma unroll @@ -95,7 +123,7 @@ T pow(T a) { } template -inline FF_CUDEV +inline FF_CUHOSTDEV T pow(T a, int N) { T p = a; # pragma unroll @@ -105,28 +133,28 @@ T pow(T a, int N) { } template -inline FF_CUDEV +inline FF_CUHOSTDEV T min(T a, T b) { return (a < b ? a : b); } template -inline FF_CUDEV +inline FF_CUHOSTDEV T max(T a, T b) { return (a > b ? a : b); } template -inline FF_CUDEV +inline FF_CUHOSTDEV T abs(T a) { return static_cast(a < 0 ? -a : a); } template -inline FF_CUDEV +inline FF_CUHOSTDEV signed char sign(T a) { return static_cast(a == 0 ? 0 : a < 0 ? -1 : 1); @@ -134,7 +162,7 @@ signed char sign(T a) #ifdef __CUDACC__ template <> -inline FF_CUDEV +inline FF_CUHOSTDEV half min<>(half a, half b) { // Compare via float: half has multiple implicit conversions to built-in @@ -144,7 +172,7 @@ half min<>(half a, half b) return (af < bf ? a : b); } template <> -inline FF_CUDEV +inline FF_CUHOSTDEV half max<>(half a, half b) { float af = static_cast(a); @@ -159,7 +187,7 @@ template ::value > struct _mod { - inline FF_CUDEV static + inline FF_CUHOSTDEV static T f(T x, U d) { signed char sx = sign(x); @@ -174,7 +202,7 @@ struct _mod template struct _mod { - inline FF_CUDEV static + inline FF_CUHOSTDEV static T f(T x, U d) { return x % d; @@ -182,14 +210,14 @@ struct _mod }; template -inline FF_CUDEV +inline FF_CUHOSTDEV T mod(T x, U d) { return _mod::f(x, d); } template -inline FF_CUDEV +inline FF_CUHOSTDEV OT typed_prod(const IT * x, size_t size) { if (size == 0) @@ -201,7 +229,7 @@ OT typed_prod(const IT * x, size_t size) } template -inline FF_CUDEV +inline FF_CUHOSTDEV OT typed_prod(const IT * x) { if (size == 0) @@ -214,21 +242,21 @@ OT typed_prod(const IT * x) } template -inline FF_CUDEV +inline FF_CUHOSTDEV T prod(const T * x, size_t size) { return typed_prod(x, size); } template -inline FF_CUDEV +inline FF_CUHOSTDEV T prod(const T * x) { return typed_prod(x); } template -inline FF_CUDEV +inline FF_CUHOSTDEV void fillfrom(U out[N], const V * inp) { # pragma unroll @@ -237,7 +265,7 @@ void fillfrom(U out[N], const V * inp) } template -inline FF_CUDEV +inline FF_CUHOSTDEV void fillfrom(U out[N], const V * inp, W stride) { # pragma unroll @@ -246,7 +274,7 @@ void fillfrom(U out[N], const V * inp, W stride) } template -inline FF_CUDEV +inline FF_CUHOSTDEV void fillfrom(int N, U out[], const V * inp) { for (int n=0; n < N; ++ n) @@ -254,7 +282,7 @@ void fillfrom(int N, U out[], const V * inp) } template -inline FF_CUDEV +inline FF_CUHOSTDEV void fillfrom(int N, U out[], const V * inp, W stride) { for (int n=0; n < N; ++n, inp += stride) @@ -262,7 +290,7 @@ void fillfrom(int N, U out[], const V * inp, W stride) } template -inline FF_CUDEV +inline FF_CUHOSTDEV void fill(U * out, V inp) { auto val = static_cast(inp); @@ -272,7 +300,7 @@ void fill(U * out, V inp) } template -inline FF_CUDEV +inline FF_CUHOSTDEV void fill(U * out, V inp, W stride) { auto val = static_cast(inp); diff --git a/tests/impl-cuda/compile_probe_hostdev.cu b/tests/impl-cuda/compile_probe_hostdev.cu new file mode 100644 index 0000000..b205465 --- /dev/null +++ b/tests/impl-cuda/compile_probe_hostdev.cu @@ -0,0 +1,122 @@ +// Compile-only probe: every helper in impl/kernels/utils.h must stay +// host-callable (FF_CUHOSTDEV), and this TU is what makes nvcc say so. +// +// Why this exists (fastfields-lib#150) +// --------------------------------------------------------------------------- +// `typed_prod` and `prod` were `inline FF_CUDEV`, i.e. `__device__` only, while +// three separate host-side callers used them: +// +// * `utils.h`'s own `FF_CUHOST canUse32BitIndexMath` -> `typed_prod`, which +// every `FF_CANUSE32BITS` in src/lib-cuda goes through; +// * every `FF_CUHOST` launcher in impl/cuda/{reg_field,reg_flow, +// distance_euclidean,distance_l1,distance_mesh}.h -> `prod(size, n)`, +// to size its grid; +// * `impl/kernels/distance/mesh.h`'s `FF_CUHOST build_tree` -> `max`. +// +// nvcc did not reject any of them, and the reason is the whole point of this +// file: **nvcc diagnoses a host->device call only when the calling function is +// not itself a template.** All three callers above are templates, so the check +// was skipped, and cudafe++ emitted into the *host* object, in place of the +// callee's real body, +// +// {int volatile ___ = 1; (void)args; ::exit(___);} +// +// Everything compiled, linked, passed `-Wl,--no-undefined` and `ldd -r`, and +// would have terminated the calling process with status 1 on the first call. +// At -O1 and above (which is what CI builds CUDA with) the host compiler also +// deleted every statement after the call, `exit` being `noreturn`, so the +// dispatch downstream of it vanished from the object entirely. +// +// What this file does +// --------------------------------------------------------------------------- +// It calls each helper from a **plain, non-template `__host__` function** -- +// the one shape nvcc *does* check. Mark any of them `FF_CUDEV` again and this +// TU fails with nvcc's own +// +// error: calling a __device__ function("ff::cuda::prod<...>") from a +// __host__ function("...") is not allowed +// +// which is exactly the diagnostic the real code could not produce. It is +// compiled, never linked, never run, and costs a couple of seconds. +// +// The second half explicitly instantiates one representative `FF_CUHOST` +// launcher per impl/cuda header that calls into these helpers. An explicit +// instantiation type-checks the launcher body in nvcc's device pass, where the +// host->device call edge *is* reported -- so this half catches a bad edge +// introduced anywhere in a launcher, not just the helpers enumerated above. +// Note that compile_probe_mesh.cu documents this same technique as a "spurious +// artifact of the probe technique, not a defect". That reading was wrong: the +// error it saw was #150, reported correctly by the compiler and dismissed. +// +// What it does NOT buy: anything about device code, and anything semantic. A +// helper that is host-callable but computes the wrong thing compiles happily. + +#include +#include +#include + +namespace U = ff::cuda; + +// --------------------------------------------------------------------------- +// utils.h, one non-template host caller per helper. +// --------------------------------------------------------------------------- + +double ff_probe_hostdev_scalars(double a, double b, int n) +{ + double x = a, y = b; + U::swap(x, y); + double acc = 0; + acc += U::square(x); + acc += U::sqrt(y < 0 ? -y : y); + acc += U::pow<3>(x); + acc += U::pow(x, n); + acc += U::min(x, y); + acc += U::max(x, y); + acc += U::abs(x); + acc += U::sign(x); + acc += U::mod(x, y); + return acc; +} + +long ff_probe_hostdev_arrays(const long * size, unsigned long ndim, + long * out, const long * inp) +{ + long acc = 0; + + // The two `typed_prod` overloads and the two `prod` overloads. The dynamic + // `typed_prod` is the exact edge `canUse32BitIndexMath` takes. + acc += U::typed_prod(size, ndim); + acc += U::typed_prod(size); + acc += U::prod(size, ndim); + acc += U::prod<3>(size); + + // fill / fillfrom, static and dynamic rank, with and without a stride. + long buf[4]; + U::fillfrom<4>(buf, inp); + U::fillfrom<4>(buf, inp, 2L); + U::fillfrom(4, out, inp); + U::fillfrom(4, out, inp, 2L); + U::fill<4>(out, 1L); + U::fill<4>(out, 1L, 2L); + + acc += buf[0]; + return acc; +} + +// `canUse32BitIndexMath` itself: a host-only function template, so nvcc cannot +// check *its* body -- but instantiating it from here at least keeps the +// signature honest, and it is the call every src/lib-cuda dispatch makes. +bool ff_probe_hostdev_canuse32(int ndim, const long * size, const long * stride) +{ + return U::canUse32BitIndexMath(ndim, size, stride); +} + +// --------------------------------------------------------------------------- +// One explicit instantiation per launcher header that reaches those helpers. +// This is the half that generalises: it re-checks the whole launcher body. +// --------------------------------------------------------------------------- + +template void ff::cuda::distance_e::dt( + int, float *, float, const int *, const int *, intptr_t); +template void ff::cuda::distance_l1::dt( + int, float *, float, const int *, const int *, intptr_t); diff --git a/tests/impl-cuda/compile_probe_mesh.cu b/tests/impl-cuda/compile_probe_mesh.cu index 0408fa1..031c1df 100644 --- a/tests/impl-cuda/compile_probe_mesh.cu +++ b/tests/impl-cuda/compile_probe_mesh.cu @@ -45,15 +45,34 @@ // compiles perfectly and silently under-launches. That class needs review or // real hardware. // -// IMPORTANT -- instantiate by *calling* from a __host__ function, not with an -// explicit-instantiation definition (`template void ...sdt<...>(...);`). An -// explicit instantiation also instantiates the body in nvcc's *device* pass, -// which reports a spurious -// error: calling a __device__ function("ff::cuda::prod") from a -// __host__ function("sdt") is not allowed -// because `prod` is FF_CUDEV. `distance_euclidean.h`'s `dt` reproduces that -// identically yet compiles fine in the real build, so it is an artifact of the -// probe technique, not a defect. See fastfields-cuda-impl#40. +// HISTORICAL NOTE -- this file used to carry the following warning: +// +// IMPORTANT -- instantiate by *calling* from a __host__ function, not with +// an explicit-instantiation definition (`template void ...sdt<...>(...);`). +// An explicit instantiation also instantiates the body in nvcc's *device* +// pass, which reports a spurious +// error: calling a __device__ function("ff::cuda::prod") from a +// __host__ function("sdt") is not allowed +// because `prod` is FF_CUDEV. `distance_euclidean.h`'s `dt` reproduces +// that identically yet compiles fine in the real build, so it is an +// artifact of the probe technique, not a defect. +// +// **That error was not spurious. It was fastfields-lib#150**, reported +// correctly by the compiler and written off. `prod` really was `__device__` +// only, every FF_CUHOST launcher in impl/cuda/ really did call it from host +// code, and the reason the real build "compiled fine" is that nvcc only checks +// host->device calls when the calling function is not itself a template: the +// launchers are templates, so instead of an error, cudafe++ silently replaced +// `prod`'s body in the host object with `::exit(1)` and the host compiler +// deleted everything after the call. The explicit instantiation was the only +// thing in the tree that made nvcc say so. +// +// `prod` is FF_CUHOSTDEV now, and tests/impl-cuda/compile_probe_hostdev.cu +// keeps a small set of explicit instantiations precisely *because* that +// technique surfaces this class of defect. The advice above is inverted where +// this call edge is concerned; what remains true is that the two techniques +// check different passes, and this file's call-based probes are what type-check +// the launch arguments. #include diff --git a/tools/check-cuda-host-stubs.sh b/tools/check-cuda-host-stubs.sh new file mode 100755 index 0000000..77c7989 --- /dev/null +++ b/tools/check-cuda-host-stubs.sh @@ -0,0 +1,68 @@ +#!/bin/sh +# Fail if any nvcc-produced object references `exit` -- the signature of a +# __host__ function calling a __device__-only one (fastfields-lib#150). +# +# WHAT IS BEING DETECTED +# --------------------------------------------------------------------------- +# nvcc rejects a host->device call only when the CALLING function is not itself +# a template. Every host caller in this tree is a template (the FF_CUHOST +# launchers in impl/cuda/, `canUse32BitIndexMath`, mesh.h's `build_tree`), so +# the check is skipped and cudafe++ instead emits, into the HOST object, this +# body in place of the callee's real one: +# +# {int volatile ___ = 1; (void)args; ::exit(___);} +# +# That compiles, links, and passes both `-Wl,--no-undefined` and `ldd -r`, +# because the damage is intra-TU: nothing becomes undefined. It terminates the +# calling process with status 1 at the first call, and -- since `exit` is +# `noreturn` -- at -O1 and above the host compiler deletes every statement +# after it, so the dtype/dim/bound dispatch downstream vanishes from the object +# too. `libfastfields-cuda.so` shipped in that state for months with every job +# green. +# +# WHY `exit` IS A SOUND ORACLE HERE +# --------------------------------------------------------------------------- +# Nothing in this codebase calls `::exit` (or `std::exit`) -- checked; the only +# `exit` in the tree is `ff::cpu::ThreadPool::Worker::exit()`, a member +# function with a different symbol. So an undefined `exit` in a lib-cuda object +# has exactly one source: a cudafe++ stub. If a deliberate `::exit` call is +# ever added below `impl/`, this check has to be revisited rather than +# silenced. +# +# This is the catch-all net. The precise, earlier gate is +# tests/impl-cuda/compile_probe_hostdev.cu, which turns the same mistake into +# an nvcc error at compile time by calling each helper from a non-template host +# function -- the one shape nvcc does diagnose. +# +# Usage: tools/check-cuda-host-stubs.sh [ ...] + +set -eu + +[ $# -gt 0 ] || { echo "usage: $0 [...]" >&2; exit 2; } + +status=0 +checked=0 +for obj in "$@"; do + [ -f "$obj" ] || { echo "no such object: $obj" >&2; status=1; continue; } + checked=$((checked + 1)) + if nm -u "$obj" 2>/dev/null | grep -qw 'exit'; then + echo "FAIL $obj -- undefined reference to exit()" + echo " A __host__ function in this TU calls a __device__-only one." + echo " cudafe++ replaced the callee's host body with exit(1); every" + echo " statement after the call was then dropped as unreachable." + echo " Find it with:" + echo " nvcc ... --keep --keep-dir /tmp/keep # then grep the" + echo " # generated .cudafe1.cpp for '::exit(___)' and read the" + echo " # function above it -- that is the mis-qualified callee." + echo " Fix: give that function FF_CUHOSTDEV, not FF_CUDEV, and add" + echo " it to tests/impl-cuda/compile_probe_hostdev.cu." + status=1 + else + echo "ok $obj" + fi +done + +if [ "$status" -eq 0 ]; then + echo "no host-pass exit(1) stubs in $checked object(s)" +fi +exit "$status"