libfastfields-cuda.so: every dispatch after FF_CANUSE32BITS is silently deleted by nvcc (__host__ calling a __device__-only typed_prod)

[balbasty]

Found while measuring the reg_flow TU split (#147). It is not caused by that change — it is pre-existing on main and affects the shipped CUDA library.

Claim

In libfastfields-cuda.so, every public entry point that computes
use_32bits = FF_CANUSE32BITS(...) and then dispatches has had the dispatch removed by the compiler. The function validates its arguments and returns without launching anything.

That is all 11 modules in src/lib-cuda/, 78 FF_CANUSE32BITS sites:

distance 19 posdef 10 pushpull 4 pushpull_backward 11 reg_field 11
reg_field_rls 4 reg_flow 12 reg_flow_rls 4 resize 1 restrict 1 splinc 1

src/lib-cpu uses the same macro and is unaffected: under the host compiler FF_CUDEV expands to nothing. That is precisely why the 59,886-check CPU gate has never seen this, and why no CUDA job has either — build-cuda is compile+link only and there is no GPU in CI.

Root cause

include/fastfields/impl/kernels/utils.h:

• line 192-194 — the dynamic-size typed_prod is inline FF_CUDEV, i.e. __device__ only:

template <typename OT, typename IT, typename size_t>
inline FF_CUDEV
OT typed_prod(const IT * x, size_t size)

• line 667-668 — canUse32BitIndexMath is FF_CUHOST inline (__host__) and calls it:

FF_CUHOST inline bool canUse32BitIndexMath(...)
{
 ...
 int64_t numel = typed_prod<int64_t>(size, ndim);

A __host__ function calling a __device__-only function template. nvcc does not reject it, and at -O1/-O2/-O3 the host object silently loses everything in the caller after that call.

prod(const T * x, size_t size) (~line 217) is FF_CUDEV too and forwards to the same typed_prod, so it is in the same class.

Minimal reproducer

Self-contained, ~40 lines, no project headers needed beyond the qualifier macros. refs counts undefined references to the three dispatch arms; 3 is correct, 0 means the switch was deleted.

#include <stdexcept>
#include <cstdint>
#include <limits>

#ifdef BUG
# define PROD_QUAL __device__ // as typed_prod is today
#else
# define PROD_QUAL __host__ __device__ // the fix
#endif

template <typename OT, typename IT>
PROD_QUAL inline OT my_prod(const IT * p, int32_t n)
{ OT r = 1; for (int32_t i = 0; i < n; ++i) r *= (OT)p[i]; return r; }

template <class N, class S, class T>
__host__ inline bool my_canUse(N ndim, const S * size, const T * stride)
{
 int64_t max32 = std::numeric_limits<int32_t>::max();
 int64_t numel = my_prod<int64_t>(size, ndim);
 if (numel >= max32) return false;
 if (numel == 0) return max32 > 0;
 if (stride == nullptr) return true;
 int64_t offset = 0, lin = numel - 1;
 for (N i = ndim - 1; i >= 0; --i) { offset += (lin % size[i]) * stride[i]; lin /= size[i]; }
 return offset < max32;
}

struct Tensor { int32_t ndim; int64_t * shape; int64_t * strides; };
namespace ns {
__attribute__((visibility("hidden"))) void f1(Tensor&, bool);
__attribute__((visibility("hidden"))) void f2(Tensor&, bool);
__attribute__((visibility("hidden"))) void f3(Tensor&, bool);
}
void entry(Tensor & t, int ndim)
{
 const bool u = my_canUse(t.ndim, t.shape, t.strides);
 switch (ndim) {
 case 1: return ns::f1(t, u);
 case 2: return ns::f2(t, u);
 case 3: return ns::f3(t, u);
 default: throw std::invalid_argument("Only 1D, 2D and 3D");
 }
}

$ for O in 0 1 2 3; do for D in "" "-DBUG"; do
 printf "%-6s -O%s refs: " "${D:-ok}" $O
 nvcc -std=c++14 -O$O $D -x cu -Xcompiler -fPIC -c -o /tmp/r.o repro.cu 2>/dev/null
 nm -u /tmp/r.o | c++filt | grep -c "ns::f"
 done; done
ok -O0 refs: 3
-DBUG -O0 refs: 3
ok -O1 refs: 3
-DBUG -O1 refs: 0 <-- dispatch deleted
ok -O2 refs: 3
-DBUG -O2 refs: 0
ok -O3 refs: 3
-DBUG -O3 refs: 0

g++ compiling the same source directly keeps all 3 at every -O. It is the -x cu path.

nvcc 12.0.140 (nvidia-cuda-toolkit from Ubuntu apt — the exact package build-cuda installs).

Confirmation on real, unmodified code

src/lib-cuda/splinc.cpp at 85fdac7, compiled with the shipping flags and CI's -O1. DISPATCH_SPLINC computes use_32bits = FF_CANUSE32BITS(inp_out) and then switches on dtype/npoles/bound. Checking for the string literals of the throws on each side of that line:

string literal	position	shipped	use_32bits = false	typed_prod fixed
Unsupported spline order	before	1	1	1
Unsupported npoles	after	0	1	1
only floating point ...	after	0	1	1

Everything after the FF_CANUSE32BITS line is gone from splinc.o, and comes back when either that one expression is neutralised or typed_prod is made host+device. The 360 ff::cuda::splinc:: template instantiations are still in the object — front-end instantiation happens regardless — so the object stays ~8 MB and the build looks entirely healthy.

Why nothing caught it

• --no-undefined and the ldd -r step both pass: the deleted calls were intra-TU, so nothing became undefined.
• build-cuda is compile+link only.
• The per-module FFMEM budget is unaffected — the instantiations are still there.
• The CPU suite cannot see it: FF_CUDEV is empty under the host compiler.

Fix

One word, verified on the reproducer and on real splinc.cpp:

 template <typename OT, typename IT, typename size_t>
-inline FF_CUDEV
+inline FF_CUHOSTDEV
 OT typed_prod(const IT * x, size_t size)

prod(const T *, size_t) should very likely go with it, and the rest of impl/kernels/ is worth auditing for other FF_CUHOST → FF_CUDEV call edges — that direction is the dangerous one, since nvcc does not diagnose it here.

Two things worth deciding alongside the fix, because the fix alone leaves the same trap set:

1. A gate. The failure is invisible to every current job. The cheapest real check is a compile-time one: assert that a chosen ff::cuda:: entry point still references its dispatch (e.g. nm -u on one object, or a linker check against a stub). Anything that would have gone red here.
2. -Werror-style handling for nvcc's host/device call diagnostics, if it can be made to emit one for this case at all.

Scope note

I have deliberately not included this in #147 — it is a correctness fix in impl/kernels/, which triggers the whole CI matrix and overlaps #145/#146, and it deserves its own review rather than riding along with a compile-time perf change. #147's measurements are unaffected either way: the slices instantiate exactly what they are meant to instantiate, and the peak-RSS and CPU figures are the same with or without the fix.

[balbasty]

Two things worth adding, now that #143 has merged and I have re-measured on b0463bc.

FF_INDEX32=0 accidentally works around this

#143 defines, in core/dispatch.h:

#if FF_INDEX32
# define FF_CANUSE32BITS(x) (canUse32BitIndexMath(x.ndim, x.shape, x.strides))
#else
# define FF_CANUSE32BITS(x) (false)
#endif

With the axis off, canUse32BitIndexMath is never called, so the __host__ →
__device__-only typed_prod edge never occurs and there is nothing for nvcc to
mis-handle. Measured on unmodified src/lib-cuda/splinc.cpp at b0463bc, shipping
flags, -O1 — presence of the two throw string literals that sit after the
FF_CANUSE32BITS line in DISPATCH_SPLINC:

	Unsupported npoles	only floating point …
FF_INDEX32=1 (shipping default)	0	0
FF_INDEX32=0	1	1

So of the two build-cuda legs #143 added, index64 produces a functionally
correct library and index32 — the default, the one that ships — does not.

That is worth knowing for two reasons:

1. It is independent confirmation of the root cause. The symptom tracks the
   presence of the call, not anything else about the module.
2. It is a mitigation available today, with no source change:
   make cuda INDEXFLAGS="-DFF_INDEX32=0". It costs the narrow-index fast path
   (which build: make the 32-bit index axis a per-backend option (FF_INDEX32), default unchanged #143's own comment notes has never been benchmarked here, because there
   is no GPU in CI) and it happens to be the cheaper leg to build anyway. That is
   a stopgap, not the fix — the typed_prod qualifier is still wrong and would
   bite again the moment anything else calls it from host code.

The gate this needs

Repeating the point from the issue body because #143 makes it cheap: the failure
is invisible to every current job, and the two-leg matrix now gives a ready-made
oracle. The index64 leg's spline_coeff object contains the dispatch; the
index32 leg's does not. A three-line step comparing one such string, or an
nm -u check on one object, would have gone red on main months ago and would
go red again on any recurrence.

Scope, unchanged

Still deliberately not fixed in #147. That PR's measurements are unaffected —
the slices instantiate exactly what they are meant to instantiate, and the
peak-RSS/CPU figures are identical with or without the fix, since front-end
instantiation happens regardless of what the optimiser later deletes.

---

Generated by Claude Code

[balbasty]

Confirmed, fixed in #155 — with three corrections to this issue, because the
reproducer's framing understates it in one direction and the follow-up comment
overstates the mitigation in the other.

The mechanism, since the issue does not name it

nvcc does diagnose a __host__ → __device__ call. It only skips the check
when the calling function is itself a template. Measured, nvcc 12.0.140:

caller	callee	result
plain __host__ fn	__device__ fn	error
plain __host__ fn	__device__ fn template	error
__host__ fn template	__device__ fn	silent
__host__ fn template	__device__ fn template	silent

Every host caller in this tree is a template, which is why build-cuda was
green. In the silent cases cudafe++ writes, into the host object, in place
of the callee's real body (kept beside it under #if 0):

{int volatile ___ = 1;(void)p;(void)n;::exit(___);}

--Werror all-warnings and -Xcudafe --display_error_number change nothing —
there is no diagnostic to promote.

Correction 1 — -O0 is not the safe end

refs: 3 at -O0 reads as "not affected". It is affected identically; only the
object-code symptom is optimisation-dependent. Linked and run (no GPU needed):

	-O0	-O1
host+device prod	f1 reached / returned normally, exit 0	same
__device__-only prod	no output, exit 1	no output, exit 1

The dispatch is still in the -O0 object and still never reached — exit(1)
runs first. This is process termination at every -O, not dead code at -O1+.

Correction 2 — FF_INDEX32=0 is not a workaround, and the index64 leg was not correct either

canUse32BitIndexMath is one of three host callers of these helpers, and
only it is behind FF_INDEX32. The other two:

• 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.

Unmodified b0463bc, src/lib-cuda/distance.cpp, -O1, FF_INDEX32=0:

exported entry points = 6 reaching exit(1) = 3
 ff::cuda::dt_euclidean ff::cuda::dt_l1 ff::cuda::dt_mesh

Exactly the three whose launchers live in those headers; the spline ones, which
have no host prod, are clean. Turning the axis off moves the truncation from
the entry point down into the launcher. It does not remove it. So the
make cuda INDEXFLAGS="-DFF_INDEX32=0" stopgap should not be relied on, and
index64 was not producing a functionally correct library.

CI corroborates: on #155 the only module whose peak RSS moves in the index64
leg is distance, +15.2% — the work that leg was discarding.

Correction 3 — the blast radius is the whole public API

"78 sites" counts FF_CANUSE32BITS lines (there are 123 occurrences), and that
is not the right unit: one entry point can hold several, and the launcher edge
is not counted at all. Walking each object's call graph to the exit stub, at
b0463bc:

module	exported ff::cuda:: entry points	reaching exit(1)
distance	6	6
posdef	8	8
pushpull	4	4
pushpull_backward	4	4
reg_field	13	13
reg_field_rls	3	3
reg_flow	13	13
reg_flow_rls	3	3
resize / restrict / splinc	1 / 1 / 1	1 / 1 / 1
total	57	57

57 of 57. On #155: 0 of 57.

The compiler had already reported this, and it was written off

tests/impl-cuda/compile_probe_mesh.cu carries, and has carried since the mesh
port:

An explicit instantiation … reports a spurious error: calling a __device__ function("ff::cuda::prod") from a __host__ function("sdt") is not allowed …
it is an artifact of the probe technique, not a defect.

It was not an artifact. It was this bug, diagnosed correctly, in the one place
in the tree where the check could fire. #155 corrects that comment and turns the
technique into a gate.

Fix and gates (#155)

All 28 FF_CUDEV in the first half of impl/kernels/utils.h → FF_CUHOSTDEV
(the second half already was; nothing in the header is device-specific), plus:

• tests/impl-cuda/compile_probe_hostdev.cu — calls each helper from a plain,
  non-template __host__ function, the shape nvcc does check. 22 errors
  against the old header, 2.5 s against the new one. Runs in compile-probe-cuda.
• tools/check-cuda-host-stubs.sh — build-cuda step, fails on any undefined
  exit in build/obj/lib-cuda/*.o. 2 s, catch-all.

One more thing worth recording

The FF_MEM_BUDGET_KB gate could not have caught this and cannot confirm the
repair. FFMEM is ru_maxrss over the nvcc tree; the restored work is all in
the host pass, and for the heavy modules the peak is cicc/ptxas in the
device pass. reg_flow moves +0.01% (13,604,176 → 13,605,772 kB) — the budget
does not trip and the headroom is unchanged. The real cost shows up where the
host pass is the peak (distance +23.5%, posdef +11.2%, the .so link +23.3%)
and in compile time (reg_field +8.6%, pushpull_backward +17.7%; whole compile
step 34m24s → 35m45s, +3.9%).

No runtime claims — there is still no GPU in CI. All #155 establishes is that
the code is present and reachable.

---

Generated by Claude Code