From 4785bc1bad641465073656dbc891991d7957bc8b Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 19 Aug 2026 19:17:11 +0000 Subject: [PATCH 1/5] fix(cuda): compile posdef/resize/restrict/splinc, and make the hub link 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. --- .github/workflows/ci.yml | 72 +++++++++++++++++++++++++++++++++++----- CLAUDE.md | 16 ++++++++- make/common.mk | 33 +++++++++++++----- src/lib-cuda/Makefile | 13 ++++++++ src/lib/Makefile | 13 +++++++- 5 files changed, 128 insertions(+), 19 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 787a8ac..870c694 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -479,13 +479,14 @@ jobs: # 16 GB runner under -j2; two ~7 GB ones do not. That split is a memory # measure -- do not recombine it to tidy the file. # - # KNOWN GAP, tracked as fastfields-lib#80 and NOT addressed here: that same - # MODULES list omits posdef, resize, restrict and splinc even though the - # .cpp files exist, so eleven FF_CUDA:: symbols the hub calls unconditionally - # are undefined in libfastfields-cuda.so. Nothing catches it because the hub - # links without -Wl,--no-undefined. This job going green therefore does NOT - # mean the CUDA library is complete. Adding those four modules changes nvcc - # peak memory and so is a measured change for its own PR, not a drive-by. + # This job also *links the hub against the CUDA backend*, which is the only + # place -Wl,--no-undefined (make/common.mk's $(NO_UNDEFINED)) actually gets + # to fire on the CUDA side. Building libfastfields-cuda.so alone cannot + # detect a module missing from MODULES: nothing inside the CUDA library + # references those entry points -- the hub does. That was fastfields-lib#80, + # where posdef/resize/restrict/splinc were absent from MODULES and eleven + # FF_CUDA:: symbols the hub calls unconditionally were undefined, with this + # job green throughout. Do not reduce this job back to `make cuda`. build-cuda: name: build-cuda (compile + link) needs: changes @@ -497,13 +498,66 @@ jobs: - name: Install CUDA toolkit and clang # `nvidia-cuda-toolkit` from apt installs CUDA ~12.x (nvcc), which # drives the build; clang is the host compiler nvcc hands off to. - run: sudo apt-get update && sudo apt-get install -y nvidia-cuda-toolkit clang + # GNU `time` is not in the runner image and is what the memory probe + # below uses; it is a few hundred kB. + run: sudo apt-get update && sudo apt-get install -y nvidia-cuda-toolkit clang time - name: Compile and link libfastfields-cuda.so # BOUNDFLAGS/SPLINEFLAGS are deliberately NOT passed here: they live # outside CXXFLAGS in src/lib-cuda/Makefile precisely so that this # -O1 override cannot silently drop the bound/spline policy, and the # Makefile's own defaults are the shipping policy. - run: make -C . cuda -j2 CXX=clang++ CXXFLAGS="-std=c++14 -O1" + # + # nvcc is wrapped in `/usr/bin/time` so every module's peak RSS is on + # the record. The MODULES split and the -j2 ceiling above are stated + # in the source as measured numbers; measuring them on every CUDA run + # is what keeps those numbers from quietly going stale as the template + # matrix grows. `time` reports the max RSS over nvcc and the children + # it waits for, so the figure is the peak of the heaviest process in + # the tree -- in practice cicc or ptxas, not the nvcc driver. + # NVCC is `?=` in make/common.mk, so this override takes effect; the + # inner single quotes are consumed by the recipe's shell, not by make. + run: | + set -o pipefail + make -C . cuda -j2 CXX=clang++ CXXFLAGS="-std=c++14 -O1" \ + NVCC="/usr/bin/time -f 'FFMEM %M kB %e s %C' nvcc" \ + 2>&1 | tee /tmp/build-cuda.log + - name: Peak nvcc memory per module + if: always() + # One line per compile, biggest first. Read against the ~3.8 GB + # per-module figure recorded in src/lib-cuda/Makefile: two concurrent + # jobs must fit a 16 GB runner. + run: | + grep -h '^FFMEM ' /tmp/build-cuda.log \ + | sed -E 's#^FFMEM ([0-9]+) kB +([0-9.]+) s .*-o ([^ ]+).*#\1 kB \2 s \3#' \ + | sort -rn | head -40 + - 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 + # src/lib-cuda's MODULES becomes a link error. Implies `make cpu` + # (the hub links -lfastfields-cpu too), which is why clang is + # installed above. + # + # CXXFLAGS is deliberately NOT overridden here. src/lib/Makefile adds + # -DFF_WITH_CUDA via `CXXFLAGS +=`, and a command-line CXXFLAGS wins + # over `+=` wholesale in GNU make -- an -O1 override would silently + # drop the define and compile the hub with no CUDA dispatch at all, + # making this step link successfully while testing nothing. The CPU + # and hub objects are the same ones the test-cpu legs build at -O3. + run: make -C . lib -j2 USE_CUDA=1 CXX=clang++ + - name: Confirm no undefined FF_CUDA symbols remain + # Belt to --no-undefined's braces, and a readable failure if the flag + # is ever dropped: nothing in libfastfields.so should still be looking + # for an ff::cuda:: symbol after the link. + run: | + set -euo pipefail + missing=$(nm -DC --undefined-only build/libfastfields.so \ + | grep 'ff::cuda::' || true) + if [ -n "$missing" ]; then + echo "$missing" + echo "::error::libfastfields.so has undefined ff::cuda:: symbols" + exit 1 + fi + echo "no undefined ff::cuda:: symbols in build/libfastfields.so" # The CUDA impl layer's compile-only probe: tests/impl-cuda/*.cu are # compiled (not run) so that launchers with no other caller -- the mesh diff --git a/CLAUDE.md b/CLAUDE.md index e439760..cdfa400 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -111,6 +111,13 @@ the CPU path is the tested source of truth and CUDA is **compile+link only**. kernels/cpu/hub changes; `test-hub` on hub changes; `build-cuda` and `compile-probe-cuda` on kernels/cuda changes. +`build-cuda` builds `libfastfields-cuda.so` **and then links the hub against it** +(`make lib USE_CUDA=1`). That second step is the point: building the CUDA +library alone cannot notice an entry point missing from it, because nothing +inside the library references those entry points — the hub does. It also prints +each module's peak `nvcc` RSS, so the memory figures quoted in +`src/lib-cuda/Makefile` can be re-checked rather than trusted. + **A change under `impl/kernels/` (or `core/`, or the build system) triggers everything** — both backends compile them. Only `api/`-only or `src/lib`-only changes may skip CUDA. @@ -123,7 +130,14 @@ pushpull's fully-static order×bound compile is nightly - **C++11** for the CPU and hub layers; **C++14** for the CUDA layer (nvcc). Object rules need `-fPIC`. Adding a module means adding it to `MODULES` in `src/lib-cpu`, `src/lib-cuda` **and** `src/lib`. -- `libfastfields.so` links `-lfastfields-cpu` with an `$ORIGIN/../lib` rpath. +- `libfastfields.so` links `-lfastfields-cpu` with an `$ORIGIN/../lib` rpath, + and with `-Wl,--no-undefined` (`$(NO_UNDEFINED)` in `make/common.mk`, cleared + on macOS/Windows where the linker has no such option). A shared object is + otherwise allowed to carry unresolved symbols, which is how fastfields-lib#80 + hid: four modules missing from `src/lib-cuda`'s `MODULES` left eleven + `FF_CUDA::` entry points undefined with every build green. The hub link is + now the gate on backend completeness — forget a `MODULES` entry and it fails + there. - Op renames from the impl layer: `resize -> resample`, `restrict -> restriction`, `splinc -> spline_coeff` (a namespace cannot share a name with a function inside `ff::cpu`). **`restriction` accumulates into diff --git a/make/common.mk b/make/common.mk index 6e6cfd0..a2eb3d5 100644 --- a/make/common.mk +++ b/make/common.mk @@ -57,6 +57,20 @@ SONAME_PREFIX = -Wl,-$(SONAME), SONAME_FLAG = $(if $(SONAME_PREFIX),$(SONAME_PREFIX)$(@F)) RPATH = -Wl,-rpath,'$$ORIGIN'/../lib +# Refuse to produce a shared object with unresolved symbols. A shared library is +# allowed to have them by default, which is how fastfields-lib#80 stayed green: +# four modules were missing from src/lib-cuda's MODULES, so eleven FF_CUDA:: +# entry points the hub calls unconditionally were simply absent from +# libfastfields-cuda.so, and neither the CUDA compile nor the hub link said a +# word -- the failure was deferred to a GPU load that CI never performs. With +# this flag that class of omission is a link error. +# +# GNU ld / lld spelling. ld64 (macOS) rejects --no-undefined and already +# defaults to erroring for a dylib; the PE/COFF linker has no equivalent and +# likewise cannot leave symbols unresolved. So both non-Linux blocks below clear +# it, and link rules must reference the variable rather than the flag. +NO_UNDEFINED = -Wl,--no-undefined + ######################################################################## # Compiler detection ######################################################################## @@ -97,16 +111,18 @@ endif ##### macOS ##### ifeq (Darwin,$(PLATFORM)) - OMPFLAG = -fopenmp=libiomp5 - SOSUF = dylib - SONAME = install_name - RPATH = -Wl,-rpath,@loader_path/../lib + OMPFLAG = -fopenmp=libiomp5 + SOSUF = dylib + SONAME = install_name + RPATH = -Wl,-rpath,@loader_path/../lib + NO_UNDEFINED = endif ifeq (arm64,$(PLATFORM)) - OMPFLAG = -fopenmp=libiomp5 - SOSUF = dylib - SONAME = install_name - RPATH = -Wl,-rpath,@loader_path/../lib + OMPFLAG = -fopenmp=libiomp5 + SOSUF = dylib + SONAME = install_name + RPATH = -Wl,-rpath,@loader_path/../lib + NO_UNDEFINED = endif ##### Windows (native OS=Windows_NT, or a Unix-like shell: MINGW*/MSYS) ##### @@ -139,6 +155,7 @@ ifdef IS_WINDOWS PICFLAG = SONAME_PREFIX = RPATH = + NO_UNDEFINED = endif ######################################################################## diff --git a/src/lib-cuda/Makefile b/src/lib-cuda/Makefile index 67cbd54..67f516d 100644 --- a/src/lib-cuda/Makefile +++ b/src/lib-cuda/Makefile @@ -70,8 +70,21 @@ SPLINEFLAGS ?= -DFF_STATIC_SPLINES=0 \ # parallelism measure, not a correctness one -- the boundary-condition policy # above is what actually brought this build back from a ~16 GB ptxas OOM. # CI additionally passes -O1 on the command line, for the same reason. +# +# MODULES must list every .cpp in this directory. It did not until +# fastfields-lib#80: posdef, resize, restrict and splinc were present as +# sources and shipped as headers, but never compiled and never linked, so the +# eleven FF_CUDA:: entry points they define (the whole Posdef family, +# `resample`, `restriction`, `spline_coeff`) were undefined symbols in +# libfastfields-cuda.so while every build reported success. The hub's CUDA link +# now passes $(NO_UNDEFINED) (see src/lib/Makefile), so a repeat of that +# omission fails the link instead of surfacing at run time on a GPU. MODULES = \ distance \ + posdef \ + resize \ + restrict \ + splinc \ reg_field \ reg_field_rls \ reg_flow \ diff --git a/src/lib/Makefile b/src/lib/Makefile index db60a77..9b3a0bb 100644 --- a/src/lib/Makefile +++ b/src/lib/Makefile @@ -49,9 +49,20 @@ $(LIBDIR)/libfastfields-cuda.$(SOSUF): # cannot run this link step before the backend libraries it links against # actually exist. ($^ would pull the .so prerequisites into the link line too, # so the recipe lists $(OBJECTS) explicitly instead.) +# +# $(NO_UNDEFINED) (-Wl,--no-undefined on GNU ld/lld; empty where the platform +# linker has no such option) makes this link the gate on backend completeness. +# The hub dispatches to FF_CPU:: and, under FF_WITH_CUDA, FF_CUDA:: symbols it +# only declares; a shared object may carry unresolved symbols, so without this +# flag a backend that simply never compiled the module defining one of them +# links clean and fails on a GPU instead. That is exactly fastfields-lib#80: +# src/lib-cuda's MODULES omitted posdef/resize/restrict/splinc, leaving eleven +# FF_CUDA:: entry points undefined with every build green. The flag applies to +# the CPU-only link too -- libfastfields-cpu.so is complete, so it costs +# nothing and keeps the same guarantee on the leg CI runs on every push. $(TARGET): $(OBJECTS) $(CPU_DEP) $(CUDA_DEP) | $(BUILDDIR) $(CXX) $(CXXFLAGS) -shared $(PICFLAG) $(SONAME_FLAG) $(RPATH) \ - -L$(LIBDIR) -lfastfields-cpu $(CUDA_LDFLAGS) \ + $(NO_UNDEFINED) -L$(LIBDIR) -lfastfields-cpu $(CUDA_LDFLAGS) \ -o $@ $(OBJECTS) # -MMD -MP emit header dependency files (*.d). This layer had none, which in a From 64c2591c985c8a49be52c31bf2169ff05d10c510 Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 19 Aug 2026 19:24:38 +0000 Subject: [PATCH 2/5] ci: verify the hub's backend symbols with ldd -r, not nm 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. --- .github/workflows/ci.yml | 48 ++++++++++++++++++++++++++++++---------- 1 file changed, 36 insertions(+), 12 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 870c694..b7860d8 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -541,23 +541,47 @@ jobs: # -DFF_WITH_CUDA via `CXXFLAGS +=`, and a command-line CXXFLAGS wins # over `+=` wholesale in GNU make -- an -O1 override would silently # drop the define and compile the hub with no CUDA dispatch at all, - # making this step link successfully while testing nothing. The CPU - # and hub objects are the same ones the test-cpu legs build at -O3. - run: make -C . lib -j2 USE_CUDA=1 CXX=clang++ - - name: Confirm no undefined FF_CUDA symbols remain + # making this step link successfully while testing nothing. + # + # BOUNDFLAGS/SPLINEFLAGS *are* overridden, to the all-Dynamic policy. + # Those live outside CXXFLAGS precisely so they can be set without + # touching it, and the question this step asks -- is every symbol the + # hub references actually defined? -- is answered identically under + # any bound/spline policy, because the policy changes which template + # instantiations exist inside a module, never which entry points it + # exports. The all-static default would rebuild the full CPU matrix at + # -O3 for no extra signal; the Dynamic policy is the cheapest compile + # that produces the same libraries' worth of exported symbols. + # Cross-policy correctness is the test-cpu matrix's job. + run: | + make -C . lib -j2 USE_CUDA=1 CXX=clang++ \ + BOUNDFLAGS="-DFF_STATIC_BOUNDS=0" \ + SPLINEFLAGS="-DFF_STATIC_SPLINES=0" + - name: Confirm the hub resolves every backend symbol at load time # Belt to --no-undefined's braces, and a readable failure if the flag - # is ever dropped: nothing in libfastfields.so should still be looking - # for an ff::cuda:: symbol after the link. + # is ever dropped. + # + # This 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 shows up as `U` in the output's dynamic symbol + # table -- that is what a DT_NEEDED reference looks like. So an nm + # check flags every single FF_CPU:: and FF_CUDA:: call as "undefined" + # on a perfectly good build. `ldd -r` does the real relocation and + # reports only what genuinely cannot be resolved. + # + # LD_LIBRARY_PATH is required: the RPATH baked into libfastfields.so + # is $ORIGIN/../lib, which is right for the *installed* layout that + # fastfields-dlpack produces but does not point at build/lib from + # build/ in the source tree. run: | set -euo pipefail - missing=$(nm -DC --undefined-only build/libfastfields.so \ - | grep 'ff::cuda::' || true) - if [ -n "$missing" ]; then - echo "$missing" - echo "::error::libfastfields.so has undefined ff::cuda:: symbols" + out=$(LD_LIBRARY_PATH="$PWD/build/lib" ldd -r build/libfastfields.so) + echo "$out" + if printf '%s\n' "$out" | grep -E 'undefined symbol|not found'; then + echo "::error::libfastfields.so does not fully resolve against its backends" exit 1 fi - echo "no undefined ff::cuda:: symbols in build/libfastfields.so" + echo "libfastfields.so resolves cleanly against both backends" # The CUDA impl layer's compile-only probe: tests/impl-cuda/*.cu are # compiled (not run) so that launchers with no other caller -- the mesh From 38a7d44f641d5fc041aaba59933d2df36b42f151 Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 19 Aug 2026 20:29:37 +0000 Subject: [PATCH 3/5] fix(cuda): qualify std::sqrt in splinc's host-side get_poles_host 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(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::::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. --- src/lib-cpu/splinc.cpp | 18 ++++++++++++------ src/lib-cuda/splinc.cpp | 23 +++++++++++++++++------ 2 files changed, 29 insertions(+), 12 deletions(-) diff --git a/src/lib-cpu/splinc.cpp b/src/lib-cpu/splinc.cpp index f21172d..6cc1c7d 100644 --- a/src/lib-cpu/splinc.cpp +++ b/src/lib-cpu/splinc.cpp @@ -29,6 +29,12 @@ FF_NAMESPACE_BEGIN(FF_DEVICE) ***********************************************************************/ // Host-side poles / npoles (mirrors kernels/splinc.h get_poles). +// +// `std::sqrt` is qualified to stay identical to src/lib-cuda's copy, where the +// qualification is mandatory: unqualified `sqrt` there resolves to the CUDEV +// (__device__) `ff::cuda::sqrt` and nvcc refuses to call it from this __host__ +// function. Here `ff::cpu::sqrt` is an ordinary host function, so either +// spelling compiles -- keep them the same anyway. static inline int get_poles_host(int order, double * poles) { switch (order) { @@ -36,18 +42,18 @@ static inline int get_poles_host(int order, double * poles) case 1: return 0; case 2: - poles[0] = sqrt(8.) - 3.; + poles[0] = std::sqrt(8.) - 3.; return 1; case 3: - poles[0] = sqrt(3.) - 2.; + poles[0] = std::sqrt(3.) - 2.; return 1; case 4: - poles[0] = sqrt(664. - sqrt(438976.)) + sqrt(304.) - 19.; - poles[1] = sqrt(664. + sqrt(438976.)) - sqrt(304.) - 19.; + poles[0] = std::sqrt(664. - std::sqrt(438976.)) + std::sqrt(304.) - 19.; + poles[1] = std::sqrt(664. + std::sqrt(438976.)) - std::sqrt(304.) - 19.; return 2; case 5: - poles[0] = sqrt(67.5 - sqrt(4436.25)) + sqrt(26.25) - 6.5; - poles[1] = sqrt(67.5 + sqrt(4436.25)) - sqrt(26.25) - 6.5; + poles[0] = std::sqrt(67.5 - std::sqrt(4436.25)) + std::sqrt(26.25) - 6.5; + poles[1] = std::sqrt(67.5 + std::sqrt(4436.25)) - std::sqrt(26.25) - 6.5; return 2; case 6: poles[0] = -0.48829458930304475513011803888378906211227916123937760839; diff --git a/src/lib-cuda/splinc.cpp b/src/lib-cuda/splinc.cpp index 39e7382..b03d38b 100644 --- a/src/lib-cuda/splinc.cpp +++ b/src/lib-cuda/splinc.cpp @@ -29,6 +29,17 @@ FF_NAMESPACE_BEGIN(FF_DEVICE) ***********************************************************************/ // Host-side poles / npoles (mirrors kernels/splinc.h get_poles). +// +// `std::sqrt`, explicitly qualified, is load-bearing here and not a style +// choice. This function is host-only, but it is defined inside `ff::cuda`, so +// unqualified `sqrt` finds `ff::cuda::sqrt` from impl/kernels/utils.h by +// ordinary lookup in the enclosing namespace -- and under __CUDACC__ that +// overload is CUDEV (__device__), which a __host__ function may not call. +// nvcc rejected this file outright until the calls were qualified; nothing +// noticed, because the module was missing from src/lib-cuda/Makefile's MODULES +// and so had never been compiled (fastfields-lib#80). src/lib-cpu's copy is +// kept identical so the two do not drift apart again -- there `ff::cpu::sqrt` +// is an ordinary host function and either spelling compiles. static inline int get_poles_host(int order, double * poles) { switch (order) { @@ -36,18 +47,18 @@ static inline int get_poles_host(int order, double * poles) case 1: return 0; case 2: - poles[0] = sqrt(8.) - 3.; + poles[0] = std::sqrt(8.) - 3.; return 1; case 3: - poles[0] = sqrt(3.) - 2.; + poles[0] = std::sqrt(3.) - 2.; return 1; case 4: - poles[0] = sqrt(664. - sqrt(438976.)) + sqrt(304.) - 19.; - poles[1] = sqrt(664. + sqrt(438976.)) - sqrt(304.) - 19.; + poles[0] = std::sqrt(664. - std::sqrt(438976.)) + std::sqrt(304.) - 19.; + poles[1] = std::sqrt(664. + std::sqrt(438976.)) - std::sqrt(304.) - 19.; return 2; case 5: - poles[0] = sqrt(67.5 - sqrt(4436.25)) + sqrt(26.25) - 6.5; - poles[1] = sqrt(67.5 + sqrt(4436.25)) - sqrt(26.25) - 6.5; + poles[0] = std::sqrt(67.5 - std::sqrt(4436.25)) + std::sqrt(26.25) - 6.5; + poles[1] = std::sqrt(67.5 + std::sqrt(4436.25)) - std::sqrt(26.25) - 6.5; return 2; case 6: poles[0] = -0.48829458930304475513011803888378906211227916123937760839; From 9b7f65aff32de699dd92d443b74355f42c1333dd Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 19 Aug 2026 21:13:32 +0000 Subject: [PATCH 4/5] build(cuda): record the measured per-module nvcc memory; tidy splinc 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. --- .github/workflows/ci.yml | 27 ++++++++++++++++++--------- src/lib-cpu/splinc.cpp | 20 ++++++++++---------- src/lib-cuda/Makefile | 8 ++++++++ src/lib-cuda/splinc.cpp | 20 ++++++++++---------- 4 files changed, 46 insertions(+), 29 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index b7860d8..923009b 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -521,15 +521,6 @@ jobs: make -C . cuda -j2 CXX=clang++ CXXFLAGS="-std=c++14 -O1" \ NVCC="/usr/bin/time -f 'FFMEM %M kB %e s %C' nvcc" \ 2>&1 | tee /tmp/build-cuda.log - - name: Peak nvcc memory per module - if: always() - # One line per compile, biggest first. Read against the ~3.8 GB - # per-module figure recorded in src/lib-cuda/Makefile: two concurrent - # jobs must fit a 16 GB runner. - run: | - grep -h '^FFMEM ' /tmp/build-cuda.log \ - | sed -E 's#^FFMEM ([0-9]+) kB +([0-9.]+) s .*-o ([^ ]+).*#\1 kB \2 s \3#' \ - | sort -rn | head -40 - 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 @@ -582,6 +573,24 @@ jobs: exit 1 fi echo "libfastfields.so resolves cleanly against both backends" + - name: Peak nvcc memory per module + # Deliberately the LAST step, and `always()` so it also reports after a + # failed compile. It is a summary of the compile step far above, and + # putting it there buried it: the hub-link step that follows emits a + # couple of hundred lines of template warnings, so the table was no + # longer anywhere near the end of the log where anyone would look. + # + # One line per nvcc invocation, biggest first. Read against the + # ~3.8 GB per-module figure recorded in src/lib-cuda/Makefile: with + # -j2, two concurrent jobs must fit a 16 GB runner. + if: always() + run: | + table=$(grep -h '^FFMEM ' /tmp/build-cuda.log \ + | sed -E 's#^FFMEM ([0-9]+) kB +([0-9.]+) s .*-o ([^ ]+).*#\1 kB \2 s \3#' \ + | sort -rn | head -40) + echo "$table" + { echo '### Peak nvcc RSS per module'; echo; echo '```'; \ + echo "$table"; echo '```'; } >> "$GITHUB_STEP_SUMMARY" # The CUDA impl layer's compile-only probe: tests/impl-cuda/*.cu are # compiled (not run) so that launchers with no other caller -- the mesh diff --git a/src/lib-cpu/splinc.cpp b/src/lib-cpu/splinc.cpp index 6cc1c7d..e2c8f45 100644 --- a/src/lib-cpu/splinc.cpp +++ b/src/lib-cpu/splinc.cpp @@ -41,19 +41,19 @@ static inline int get_poles_host(int order, double * poles) case 0: case 1: return 0; - case 2: - poles[0] = std::sqrt(8.) - 3.; - return 1; - case 3: - poles[0] = std::sqrt(3.) - 2.; - return 1; + case 2: poles[0] = std::sqrt(8.) - 3.; return 1; + case 3: poles[0] = std::sqrt(3.) - 2.; return 1; case 4: - poles[0] = std::sqrt(664. - std::sqrt(438976.)) + std::sqrt(304.) - 19.; - poles[1] = std::sqrt(664. + std::sqrt(438976.)) - std::sqrt(304.) - 19.; + poles[0] = + std::sqrt(664. - std::sqrt(438976.)) + std::sqrt(304.) - 19.; + poles[1] = + std::sqrt(664. + std::sqrt(438976.)) - std::sqrt(304.) - 19.; return 2; case 5: - poles[0] = std::sqrt(67.5 - std::sqrt(4436.25)) + std::sqrt(26.25) - 6.5; - poles[1] = std::sqrt(67.5 + std::sqrt(4436.25)) - std::sqrt(26.25) - 6.5; + poles[0] = + std::sqrt(67.5 - std::sqrt(4436.25)) + std::sqrt(26.25) - 6.5; + poles[1] = + std::sqrt(67.5 + std::sqrt(4436.25)) - std::sqrt(26.25) - 6.5; return 2; case 6: poles[0] = -0.48829458930304475513011803888378906211227916123937760839; diff --git a/src/lib-cuda/Makefile b/src/lib-cuda/Makefile index 67f516d..5a23bcd 100644 --- a/src/lib-cuda/Makefile +++ b/src/lib-cuda/Makefile @@ -79,6 +79,14 @@ SPLINEFLAGS ?= -DFF_STATIC_SPLINES=0 \ # libfastfields-cuda.so while every build reported success. The hub's CUDA link # now passes $(NO_UNDEFINED) (see src/lib/Makefile), so a repeat of that # omission fails the link instead of surfacing at run time on a GPU. +# +# Adding them does not move the memory ceiling above. Measured on the CI runner +# under the same conditions as the figures above (nvcc, -O1, the BOUNDFLAGS and +# SPLINEFLAGS defaults in this file), the four are the *cheapest* modules here: +# resize ~2.1 GB / ~190 s, restrict ~1.4 GB / ~120 s, posdef ~0.4 GB / ~40 s, +# splinc smaller still -- against ~3.8 GB for a split regulariser, which remains +# what sets the -j2 ceiling. build-cuda re-measures and prints the whole table +# on every run, so this list can be checked rather than trusted. MODULES = \ distance \ posdef \ diff --git a/src/lib-cuda/splinc.cpp b/src/lib-cuda/splinc.cpp index b03d38b..8157e27 100644 --- a/src/lib-cuda/splinc.cpp +++ b/src/lib-cuda/splinc.cpp @@ -46,19 +46,19 @@ static inline int get_poles_host(int order, double * poles) case 0: case 1: return 0; - case 2: - poles[0] = std::sqrt(8.) - 3.; - return 1; - case 3: - poles[0] = std::sqrt(3.) - 2.; - return 1; + case 2: poles[0] = std::sqrt(8.) - 3.; return 1; + case 3: poles[0] = std::sqrt(3.) - 2.; return 1; case 4: - poles[0] = std::sqrt(664. - std::sqrt(438976.)) + std::sqrt(304.) - 19.; - poles[1] = std::sqrt(664. + std::sqrt(438976.)) - std::sqrt(304.) - 19.; + poles[0] = + std::sqrt(664. - std::sqrt(438976.)) + std::sqrt(304.) - 19.; + poles[1] = + std::sqrt(664. + std::sqrt(438976.)) - std::sqrt(304.) - 19.; return 2; case 5: - poles[0] = std::sqrt(67.5 - std::sqrt(4436.25)) + std::sqrt(26.25) - 6.5; - poles[1] = std::sqrt(67.5 + std::sqrt(4436.25)) - std::sqrt(26.25) - 6.5; + poles[0] = + std::sqrt(67.5 - std::sqrt(4436.25)) + std::sqrt(26.25) - 6.5; + poles[1] = + std::sqrt(67.5 + std::sqrt(4436.25)) - std::sqrt(26.25) - 6.5; return 2; case 6: poles[0] = -0.48829458930304475513011803888378906211227916123937760839; From a3d26eefd854ab92d668021ff31639ef01145152 Mon Sep 17 00:00:00 2001 From: Claude Date: Wed, 19 Aug 2026 23:05:45 +0000 Subject: [PATCH 5/5] docs(cuda): correct the recorded ptxas memory figures with measured ones 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. --- .github/workflows/ci.yml | 18 ++++++++----- CLAUDE.md | 15 +++++++---- src/lib-cuda/Makefile | 55 ++++++++++++++++++++++++++++++---------- 3 files changed, 63 insertions(+), 25 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 923009b..08cb2e3 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -469,15 +469,19 @@ jobs: # bound::type::Dynamic (see BOUNDFLAGS in src/lib-cuda/Makefile), and # -O1 is the belt to that braces. # - # 2. -j2, not -j$(nproc). Each regulariser module still peaks at ~4 GB in - # ptxas, so two concurrent nvcc jobs is the safe ceiling on a 16 GB - # runner. + # 2. -j2, not -j$(nproc). The regulariser modules are enormous in ptxas and + # two concurrent nvcc jobs is the ceiling on a 16 GB runner. Note the + # "~4 GB per module" this comment used to give as the reason is not the + # real number: measured, reg_flow peaks at 12.98 GB and reg_field at + # 8.93 GB, so -j2 survives because make does not happen to overlap the + # two heaviest peaks, not because two of them fit. The full measured + # table, and what it does and does not license, is in + # src/lib-cuda/Makefile above MODULES -- read it before touching -j. # # Relatedly: src/lib-cuda/Makefile's MODULES list is deliberately split - # (reg_field / reg_field_rls, reg_flow / reg_flow_rls) because ptxas peaks - # at ~3.8 GB per split module but ~6-7 GB combined. Two ~4 GB jobs fit a - # 16 GB runner under -j2; two ~7 GB ones do not. That split is a memory - # measure -- do not recombine it to tidy the file. + # (reg_field / reg_field_rls, reg_flow / reg_flow_rls). That split is a + # memory measure and is still load-bearing at the measured sizes -- do not + # recombine it to tidy the file. # # This job also *links the hub against the CUDA backend*, which is the only # place -Wl,--no-undefined (make/common.mk's $(NO_UNDEFINED)) actually gets diff --git a/CLAUDE.md b/CLAUDE.md index cdfa400..b7eef27 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -148,11 +148,16 @@ pushpull's fully-static order×bound compile is nightly `BOUNDFLAGS` / `SPLINEFLAGS`. These live **outside** `CXXFLAGS` on purpose so that a `CXXFLAGS=` override (as CUDA CI does, to force `-O1`) cannot silently drop the policy. -- **CUDA memory limits are measured, not guessed.** CUDA CI forces `-O1` and - `-j2` because `ptxas` was OOM-killed at ~16 GB on `reg_field.cpp`, and - `src/lib-cuda`'s `MODULES` split (`reg_field`/`reg_field_rls`, - `reg_flow`/`reg_flow_rls`) caps peak memory at ~3.8 GB per module against - ~6–7 GB combined. Do not recombine or "tidy" these. +- **CUDA memory limits are measured, not guessed** — and the numbers that used + to be recorded here were wrong. CUDA CI forces `-O1` and `-j2` because + `ptxas` was OOM-killed at ~16 GB on `reg_field.cpp`, and `src/lib-cuda`'s + `MODULES` split (`reg_field`/`reg_field_rls`, `reg_flow`/`reg_flow_rls`) is + what keeps that build possible. But the "~3.8 GB per module" figure this note + carried is not what the build actually does: `build-cuda` now measures every + module and **`reg_flow` peaks at 12.98 GB of a 16 GB runner**. The measured + 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. - `include/fastfields/core/dlpack.h` is vendored upstream code: do not edit it, and it is skipped by `codespell` (see `.codespellrc`). diff --git a/src/lib-cuda/Makefile b/src/lib-cuda/Makefile index 5a23bcd..5fc84ba 100644 --- a/src/lib-cuda/Makefile +++ b/src/lib-cuda/Makefile @@ -64,12 +64,44 @@ SPLINEFLAGS ?= -DFF_STATIC_SPLINES=0 \ # The regularisers are split into a core module and an `_rls` module (the # reweighted-least-squares ops) -- a CUDA-only split that lib-cpu does not need. -# Rationale: measured with the default BOUNDFLAGS, `ptxas` peaks at ~3.8 GB per -# split module but ~6-7 GB for the combined file. Two ~4 GB jobs fit a 16 GB CI -# runner under `make -j2`; two ~7 GB ones do not. The split is a memory/ -# parallelism measure, not a correctness one -- the boundary-condition policy -# above is what actually brought this build back from a ~16 GB ptxas OOM. -# CI additionally passes -O1 on the command line, for the same reason. +# It is a memory/parallelism measure, not a correctness one: the +# boundary-condition policy above is what brought this build back from a ~16 GB +# ptxas OOM, and CI additionally passes -O1 for the same reason. +# +# ~~~ MEASURED, 2026-08-19 (fastfields-lib#80) ~~~ +# The figures this comment used to quote -- "~3.8 GB per split module, ~6-7 GB +# combined" -- are WRONG, and wrong in the dangerous direction. build-cuda now +# wraps nvcc in /usr/bin/time and prints peak RSS per module; on a 16 GB +# ubuntu-latest runner, nvcc -O1 with the BOUNDFLAGS/SPLINEFLAGS defaults above: +# +# reg_flow 12.98 GB 1097 s <-- not ~3.8 +# reg_field 8.93 GB 840 s <-- not ~3.8 +# reg_field_rls 6.67 GB 759 s <-- not ~3.8 +# 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 +# +# (Peak RSS of the largest single process in the nvcc tree -- cicc or ptxas -- +# not a sum: `time` reports ru_maxrss, which is a maximum. Reproducible to +# better than 0.1% across runs.) +# +# What that means, and what it does not: +# * The SPLIT IS STILL LOAD-BEARING and must not be undone -- reg_field and +# reg_flow are the two heaviest things here even after being halved. +# * The -j2 CEILING IS NOT AS SAFE AS THE OLD NUMBERS SUGGESTED. reg_flow +# alone takes 13 of the runner's 16 GB. -j2 survives because make does not +# happen to overlap the two heaviest peaks, not because two jobs are +# known to fit. Do NOT raise -j on the strength of the old figures, and +# treat anything that makes a regulariser heavier as a real OOM risk. +# * Nothing here is caused by the four modules added below; the regularisers +# were always the hogs and are untouched by that change. +# Deciding what to do about the reg_flow headroom is out of scope for #80 -- +# it is a pre-existing condition, now measured instead of guessed. # # MODULES must list every .cpp in this directory. It did not until # fastfields-lib#80: posdef, resize, restrict and splinc were present as @@ -80,13 +112,10 @@ SPLINEFLAGS ?= -DFF_STATIC_SPLINES=0 \ # now passes $(NO_UNDEFINED) (see src/lib/Makefile), so a repeat of that # omission fails the link instead of surfacing at run time on a GPU. # -# Adding them does not move the memory ceiling above. Measured on the CI runner -# under the same conditions as the figures above (nvcc, -O1, the BOUNDFLAGS and -# SPLINEFLAGS defaults in this file), the four are the *cheapest* modules here: -# resize ~2.1 GB / ~190 s, restrict ~1.4 GB / ~120 s, posdef ~0.4 GB / ~40 s, -# splinc smaller still -- against ~3.8 GB for a split regulariser, which remains -# what sets the -j2 ceiling. build-cuda re-measures and prints the whole table -# on every run, so this list can be checked rather than trusted. +# Adding them does not move the ceiling: in the table above the four are the +# four cheapest modules in the file (resize 2.00, restrict 1.30, splinc 0.42, +# posdef 0.37 GB), together well under a single regulariser. They cost ~7 min +# of the ~31 min compile. MODULES = \ distance \ posdef \