fix(cuda): compile posdef/resize/restrict/splinc, and make the hub link strict by balbasty · Pull Request #87 · fastfields/fastfields-lib

balbasty · 2026-08-19T19:18:04Z

Closes#80.

Verified in the consolidated tree first

Everything in the issue holds on main @ 72f9798:

src/lib-cuda/Makefile's MODULES lists 7 of the 11 .cpp files in
src/lib-cuda/: posdef, resize, restrict, splinc are missing.
All four sources are present and define exactly the 11 entry points the issue
names (sym_matvec, sym_matvec_backward, sym_addmatvec_, sym_submatvec_,
sym_solve, sym_solve_, sym_invert, sym_invert_, resample,
restriction, spline_coeff).
src/lib/{posdef,resize,restrict,splinc}.cpp dispatch to all 11 under
if (IS_CUDA(...)) — real forwards, not throw-stubs.
Cross-checked the rest of the boundary: every otherFF_CUDA:: symbol the
hub calls is defined in a module that was listed, and every FF_CPU::
symbol the hub calls is defined in src/lib-cpu (whose MODULES is
complete). So the gap is exactly these four modules, CUDA-only.

Changes

src/lib-cuda/Makefile — the four modules added to MODULES.
make/common.mk — new NO_UNDEFINED variable (-Wl,--no-undefined),
cleared in the macOS and Windows blocks alongside PICFLAG/SONAME_PREFIX/
RPATH: ld64 rejects the flag and already errors for a dylib, and the
PE/COFF linker has no equivalent and cannot leave symbols unresolved either.
src/lib/Makefile — the hub link passes $(NO_UNDEFINED), on the
CPU-only link as well as the CUDA one.
src/lib-cuda/splinc.cpp (+ its src/lib-cpu mirror) — see "second
defect" below.
.github/workflows/ci.yml — the hub is now actually linked in CI, and
nvcc memory is measured.

The second defect, found by fixing the first

Adding splinc to MODULES made the build fail: src/lib-cuda/splinc.cpp
does not compile under nvcc, and never had, because it had never been
compiled. 14 errors, all one shape:

splinc.cpp(39): error: calling a __device__ function
("T1 ff::cuda::sqrt<double>(T1)") from a __host__ function
("get_poles_host") is not allowed

get_poles_host is host-only but lives inside ff::cuda, so unqualified
sqrt is found by ordinary lookup in the enclosing namespace and binds to
ff::cuda::sqrt from impl/kernels/utils.h — which under __CUDACC__ is
CUDEV, i.e. __device__. Qualifying as std::sqrt picks the host function
that was meant; values are unchanged (the non-CUDA ff::<dev>::sqrt is itself
a forward to std::sqrt). The src/lib-cpu copy gets the identical edit — it
compiles either way, but the two are meant to be mirrors and leaving them
spelled differently is how this comes back.

posdef, resize and restrict compiled clean, so splinc was the only one
of the four hiding anything.

The CI change is not incidental

build-cuda only ran make cuda. That build can never detect this class of
bug: nothing inside libfastfields-cuda.so references its own entry points —
the hub does. -Wl,--no-undefined on the hub link is only a gate if CI performs
that link, and it did not: no job in ci.yml built libfastfields.so at all.

build-cuda now builds the CUDA library (unchanged: -O1, -j2, policy left
to the Makefile), links the hub against it (make lib USE_CUDA=1), and runs
ldd -r on the result. Verified in the log:

clang++ ... -Wl,--no-undefined -L.../build/lib -lfastfields-cpu -L.../build/lib -lfastfields-cuda -o .../build/libfastfields.so
...
libfastfields.so resolves cleanly against both backends

Two traps found while validating this, either of which would have made the new
steps wrong:

The check must be ldd -r, not nm -D --undefined-only. A symbol that
--no-undefined accepted because a shared library on the link line defines
it still appears as U in the dynamic symbol table — that is what a
DT_NEEDED reference looks like. My first version grepped nm and would have
failed on a good build: a clean CPU-only make lib carries 77 such entries,
~50 of them ff::cpu::. Fixed in 64c2591; both mechanisms were then checked
against a deliberately-broken toy library (--no-undefined errors at link,
ldd -r prints undefined symbol:). LD_LIBRARY_PATH is needed because the
baked RPATH is $ORIGIN/../lib, right for the installed layout but not for
build/ → build/lib in the source tree.
CXXFLAGS must not be overridden on the hub step.src/lib/Makefile adds
-DFF_WITH_CUDA via CXXFLAGS +=, and a command-line CXXFLAGS beats +=
wholesale in GNU make — an -O1 override would compile the hub with no CUDA
dispatch and link "successfully" while testing nothing. BOUNDFLAGS/
SPLINEFLAGSare overridden (all-Dynamic): those sit outside CXXFLAGS for
exactly this reason, and the policy changes which instantiations exist inside a
module, never which entry points it exports.

Memory: measured, and the recorded numbers were wrong

nvcc is now wrapped in /usr/bin/time and the job prints peak RSS per module.
Peak RSS of the largest single process in each nvcc tree (cicc/ptxas), -O1,
shipping BOUNDFLAGS/SPLINEFLAGS, 16 GB ubuntu-latest, reproducible to
better than 0.1% across runs:

module	peak RSS	wall
`reg_flow`	12.98 GB	1097 s
`reg_field`	8.93 GB	840 s
`reg_field_rls`	6.67 GB	759 s
`resize`	2.00 GB	188 s
`reg_flow_rls`	1.90 GB	153 s
`pushpull_backward`	1.51 GB	163 s
`pushpull`	1.48 GB	153 s
`restrict`	1.30 GB	113 s
`distance`	0.77 GB	72 s
`splinc`	0.42 GB	67 s
`posdef`	0.37 GB	34 s

The four modules this PR adds are the four cheapest in the file — together
under a single regulariser, ~7 min of the ~31 min compile. The memory story for
this change is a non-event.

But the figure the tree recorded — "~3.8 GB per split module, two ~4 GB jobs fit
a 16 GB runner under -j2" — is wrong by up to 3.4×. reg_flow alone takes 13
of the runner's 16 GB. -j2 passes because make does not happen to overlap the
two heaviest peaks, not because two of them fit. That is pre-existing and not
caused by this PR (the regularisers are untouched), but it invalidates the
stated basis for -j2, so a3d26ee replaces the guess with the table in
src/lib-cuda/Makefile and points ci.yml/CLAUDE.md at it. Deciding what
to do about the reg_flow headroom is deliberately left out of this PR.

Scope

FF_CUDA::field_cg is out of scope as the issue says — src/lib/solve_field.cpp
throws rather than dispatching, so it is link-safe. Tracked by #34.

A green build-cuda means compile + link only. There is no GPU in CI;
nothing in these four modules is executed by this PR, and no claim is made about
their GPU runtime behaviour. The shared kernel math is covered by the CPU suite
(same source, both backends).

Known residual gap (not addressed here)

The hub link is exercised on CUDA-triggering changes only. A hub-only change
still never links libfastfields.so in CI, so the CPU-side --no-undefined has
no gate of its own on that path — and make / make lib, the default build
fastfields-dlpack's setup.py invokes, is still built by no CI job at all.
Worth its own issue.

CPU baseline

Run locally against this tree, and matched line-for-line by CI's
test-cpu (clang-static): 59,886 checks across 13 suites, 0 failures —
unchanged, as expected, since the only CPU source touched is a sqrt → std::sqrt
spelling with identical semantics.

Note on the 2026-08-19 21:14–22:15 UTC CI cancellations

Four jobs on run 32302780577 show as cancelled. None of them ran a line of
this diff: each died inside its apt-get step, at exactly its timeout-minutes
budget (GitHub reports a timeout kill as cancelled, not failure).

job	died in	duration	budget
`lint (clang-format)`	`Install clang-format`	10m11s	10
`test-cpu (clang-static)`	`Install build toolchain`	45m18s	45
`test-cpu (sanitize)`	`Install build toolchain`	45m17s	45
`compile-probe-cuda`	`Install CUDA toolkit and clang`	60m36s	60

apt-get update logged its last line at 21:16:52 and stalled 8 minutes to the
kill, having fallen off azure.archive.ubuntu.com onto archive.ubuntu.com.
Repo-wide, not branch-specific: PR #86's lint (clang-format) died the same way
in the same minute (21:14:16→21:24:27). build-cuda was hit too — its apt took
55 min on an earlier run — and survived only because its budget is 120. This is
the same failure mode ci.yml already documents in the test-hub timeout
comment, which is why test-hub (deliberately apt-free) was among the jobs that
passed. The mirror had recovered by 23:03.

…nk strict (#80) src/lib-cuda/Makefile's MODULES listed seven of the eleven .cpp files in that directory. posdef, resize, restrict and splinc were never compiled and never linked, so the eleven FF_CUDA:: entry points they define -- the whole Posdef family plus resample, restriction and spline_coeff -- were absent from libfastfields-cuda.so. The hub dispatches to all eleven unconditionally under `if (IS_CUDA(...))`, so a FF_WITH_CUDA build reached an undefined symbol on GPU input for all of Posdef and all of Resampling, with every build green. Two things had to be true for that to stay hidden, and both are fixed here: * The four modules are now in MODULES. * The hub links with -Wl,--no-undefined (make/common.mk's $(NO_UNDEFINED), cleared on macOS and Windows where the platform linker has no equivalent and already refuses to leave symbols unresolved). A shared object may carry undefined symbols by default, which is why ld said nothing. It applies to the CPU-only link as well: libfastfields-cpu.so is complete, so it costs nothing there and keeps the same guarantee on the leg that runs on every push. CI could not have caught it either, and could not catch a recurrence: the build-cuda job only built libfastfields-cuda.so, and nothing inside that library references its own entry points -- the hub does. So build-cuda now also links the hub against the CUDA backend (`make lib USE_CUDA=1`) and then greps the result for undefined ff::cuda:: symbols. Note it deliberately does not override CXXFLAGS for that step: src/lib/Makefile adds -DFF_WITH_CUDA via `CXXFLAGS +=`, which a command-line CXXFLAGS would drop wholesale, leaving a step that links happily while testing nothing. The CUDA memory budget is measured, not assumed: nvcc is wrapped in /usr/bin/time for the whole build and the job prints peak RSS per module, so the ~3.8 GB-per-module figure the MODULES split is justified by can be re-checked rather than trusted. -O1 and -j2 are unchanged.

The verification step added alongside the --no-undefined gate was wrong and would have failed on a correct build. `nm -D --undefined-only` lists symbols that are resolved from a shared library on the link line as `U` -- that is simply what a DT_NEEDED reference looks like -- so it reports every FF_CPU:: and FF_CUDA:: call the hub makes, whether or not anything is actually missing. Confirmed locally: a clean CPU-only `make lib` (which --no-undefined accepted) still shows 77 undefined entries, ~50 of them ff::cpu::. `ldd -r` performs the real relocation and reports only what cannot be resolved, which is the question the step means to ask. It needs LD_LIBRARY_PATH: the RPATH is $ORIGIN/../lib, correct for the installed layout fastfields-dlpack produces but not for build/ -> build/lib in the source tree. Also drop the hub step to the all-Dynamic bound/spline policy. Those variables sit outside CXXFLAGS so they can be set independently, and the policy decides which template instantiations exist inside a module, never which entry points it exports -- so the link question is answered identically, at a fraction of the compile. Locally the Dynamic build of libfastfields-cpu.so + the hub is ~4 min at -j2 cold, against an all-static -O3 build that had not finished in over 10 at -j3.

Adding splinc to src/lib-cuda's MODULES turned up a second defect underneath the first: the file does not compile under nvcc, and never had, because it had never been compiled. 14 errors, all the same shape: splinc.cpp(39): error: calling a __device__ function ("T1 ff::cuda::sqrt<double>(T1)") from a __host__ function ("get_poles_host") is not allowed `get_poles_host` is host-only but lives inside `ff::cuda`, so unqualified `sqrt` is found by ordinary lookup in the enclosing namespace and resolves to `ff::cuda::sqrt` from impl/kernels/utils.h -- which under __CUDACC__ is CUDEV, i.e. __device__. Qualifying the calls as `std::sqrt` picks the host function that was meant; the values are unchanged (the non-CUDA `ff::<dev>::sqrt` is itself a forward to `std::sqrt`). src/lib-cpu/splinc.cpp gets the identical edit. It compiles either way, since `ff::cpu::sqrt` is an ordinary host function, but the two dispatch files are meant to be mirrors and leaving them spelled differently is how this comes back. Both carry a comment saying which way round the constraint runs. posdef, resize and restrict compiled clean on the first CUDA run, so splinc was the only one of the four hiding anything.

…formatting Numbers from the build-cuda run: nvcc peak RSS at -O1 under this file's BOUNDFLAGS/SPLINEFLAGS defaults is ~2.1 GB for resize, ~1.4 GB for restrict, ~0.4 GB for posdef, and less for splinc -- all comfortably under the ~3.8 GB a split regulariser module needs, which stays the figure that sets the -j2 ceiling on a 16 GB runner. So the four modules added for fastfields-lib#80 do not change the memory story; the comment above MODULES now says so with the measurements rather than leaving the reader to assume it. The memory table moves to the end of the build-cuda job. It summarises the compile step, but sitting directly after it the table was immediately buried by the couple of hundred template warnings the hub-link step emits, so it was no longer near the end of the log where anyone would look for it. It stays `always()` so a failed compile still reports what it had reached, and it now also lands in the job summary. splinc.cpp's poles table is reflowed to what .clang-format asks for on the lines the previous commit touched (80-column wrap, and short case labels on one line). No behaviour change; both copies stay identical.

The build-cuda job now measures what it used to assert, and the assertion was wrong. Recorded across this tree was "ptxas peaks at ~3.8 GB per split module, ~6-7 GB combined", with -j2 justified as "two ~4 GB jobs fit a 16 GB runner". Measured on the runner (nvcc -O1, the shipping BOUNDFLAGS/SPLINEFLAGS, peak RSS of the largest process in each nvcc tree): reg_flow 12.98 GB 1097 s reg_field 8.93 GB 840 s reg_field_rls 6.67 GB 759 s resize 2.00 GB 188 s reg_flow_rls 1.90 GB 153 s pushpull_backward 1.51 GB 163 s pushpull 1.48 GB 153 s restrict 1.30 GB 113 s distance 0.77 GB 72 s splinc 0.42 GB 67 s posdef 0.37 GB 34 s reg_flow alone takes 13 of the runner's 16 GB -- 3.4x the figure that was supposed to make -j2 safe. The build passes, but it passes because make does not happen to schedule the two heaviest peaks together, not because two of them are known to fit. That is worth knowing before anyone raises -j, drops -O1, turns a boundary condition back to static, or adds architectures to the nvcc command line on the strength of the old numbers. Nothing here is a consequence of adding posdef/resize/restrict/splinc: those four are the four cheapest modules in the file and the regularisers, which are untouched, were always the hogs. The split is still load-bearing at the real sizes and must not be undone. What to do about the reg_flow headroom is deliberately not decided here -- it is a pre-existing condition, now measured instead of guessed. The stale figures are corrected in src/lib-cuda/Makefile (with the table), and the copies of them in ci.yml and CLAUDE.md now point at it rather than repeating a number.

Brings in #87, #90, #91 and #95. One conflict, in include/fastfields/impl/kernels/parallel.h: #91 renamed FF_NAMESPACE_BEGIN(FF) to FF_NAMESPACE_BEGIN(FF_NS) on the line this branch inserts the FF_GRAIN_SIZE block above. Resolved by keeping both -- the new block, then main's FF_NS spelling. Everything else auto-merged. #91's renames do not touch anything this branch depends on: has_atomic_add / anyAtomicAdd keep their names, FF_NS still expands to ff, and the CUDEV -> FF_CUDEV rename is confined to the CUDA half of atomic.h. FF_GRAIN_SIZE, the one macro this branch adds to an installed header, already satisfies #91's FF_-prefix rule -- `tools/rename-macros.py --check` reports "0 file(s) would change" and "include/ is clean". Re-verified on the merge result: * tools/test-baseline.sh --legs default,lib -> byte-identical to tools/test-baseline.expected. 13 suites, 59,886 checks, 0 failures. * -DFF_GRAIN_SIZE=1 -> 59,886 / 13 / 0. * -DFF_GRAIN_SIZE=1 + TSan, FF_NUM_THREADS=4, halt_on_error=1 -> 59,886 / 13 / 0, zero reports. * clone syscalls: 0 across all 13 binaries at the shipping grain size, 2 per binary at FF_GRAIN_SIZE=1. The threshold is unchanged by the merge (0 clones at n=32768, 2 at n=32769, on main and on this branch alike). * The thread-pool defects still reproduce on main at f63c7d8: the data race is deterministic under TSan (threadpool.h:148 write / :164 read) and the lost-wakeup deadlock is stochastic (8/320 trials over FF_NUM_THREADS 8/16/32/64). Both are gone on this branch: 0/320 hangs, 0 TSan reports.

claude added 5 commits August 19, 2026 19:17

balbasty merged commit 81c60fc into mainAug 19, 2026
11 checks passed

balbasty deleted the fix/cuda-modules-no-undefined branch August 19, 2026 23:44

balbasty mentioned this pull request Aug 20, 2026
The CPU suite never runs anything in parallel — turn it on, under TSan, and fix the deadlock that finds #97
Merged

Navigation Menu

Uh oh!

fix(cuda): compile posdef/resize/restrict/splinc, and make the hub link strict - #87

fix(cuda): compile posdef/resize/restrict/splinc, and make the hub link strict#87
balbasty merged 5 commits into
mainfrom
fix/cuda-modules-no-undefined

balbasty commented Aug 19, 2026•
edited
Loading

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

balbasty commented Aug 19, 2026• edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Verified in the consolidated tree first

Changes

The second defect, found by fixing the first

The CI change is not incidental

Memory: measured, and the recorded numbers were wrong

Scope

Known residual gap (not addressed here)

CPU baseline

Note on the 2026-08-19 21:14–22:15 UTC CI cancellations

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

balbasty commented Aug 19, 2026•
edited
Loading