fix(kernels): utils.h helpers must be host+device, not device-only (#150)

[balbasty]

Fixes#150. The report is right about the defect and wrong about three of its
details — one of them the part that decides whether the shipped library works.
Everything below is measured on nvcc 12.0.140, the version build-cuda
installs.

The mechanism — what nvcc actually does

nvcc does reject a __host__ → __device__ call. But only when the
calling function is not itself a template. That is the whole hazard, and it is
not documented as a limitation anywhere in the CUDA guide.

caller	callee	nvcc 12.0.140
plain __host__ function	__device__ function	error
plain __host__ function	__device__ function template	error
__host__ function template	__device__ function	accepted, no diagnostic
__host__ function template	__device__ function template	accepted, no diagnostic

Rows 1–2 give

matrix.cu(6): error: calling a __device__ function("d_plain(const int *, int)")
from a __host__ function("h") is not allowed

Rows 3–4 give rc=0 and empty stderr. --Werror all-warnings and
-Xcudafe --display_error_number add nothing: there is no diagnostic to
promote.

In the silent rows cudafe++ writes this into the host object in place of
the callee's real body — the real one is kept beside it under #if 0:

# 11 "repro.cu"
template<classOT, classIT>__attribute__((unused)) inlineOTmy_prod(constIT*p, int32_tn)
{intvolatile___=1;(void)p;(void)n;::exit(___);}
#if0
{ OTr= (1); for (int32_ti=0; i<n; ++i) { r *= ((OT)(p[i])); } returnr; }
#endif

Everything follows from that. It is a call to exit, not an error, and exit
is noreturn, so from -O1 up the host compiler deletes every statement after
the call. The buggy entry point at -O1 is five instructions:

0000000000000019 <entry(Tensor&, int)>:
19: endbr64
1d: sub $0x18,%rsp
21: movl $0x1,0xc(%rsp)
29: mov 0xc(%rsp),%edi
2d: call 32 <...> 2e: R_X86_64_PLT32 exit-0x4

Correction 1 — -O0 is not the safe end of the range

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

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

At -O0 the dispatch is still in the object and still never reached: the
process is gone before it. This is a process-terminating defect at every
optimisation level, not dead code at -O1+. For a library loaded into a Python
process, that is exit(1) under the user's interpreter.

Correction 2 — FF_INDEX32=0 is not a workaround, and index64 is not correct

The follow-up comment concludes that "index64 produces a functionally correct
library and index32 … does not". That holds for splinc and does not
generalise. canUse32BitIndexMath is only one of three host callers:

• canUse32BitIndexMath → typed_prod, behind every FF_CANUSE32BITS;
• 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.

Only the first is behind FF_INDEX32. Measured on unmodified b0463bc,
distance.cpp, -O1, FF_INDEX32=0:

distance.o 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 the three headers that call prod
from host code. The spline entry points, whose launchers do not, are fine.
Turning the axis off moves the truncation from the entry point down into the
launcher; it does not remove it. Neither build-cuda leg produced a working
library.

CI agrees: in the index64 leg, the only module whose peak RSS moves on this
branch is distance, +15.2% (564,184 → 650,204 kB) — the work that was
being discarded there, now being done.

Correction 3 — the blast radius is 100%, and "78 sites" is not the unit

"78" counts FF_CANUSE32BITSlines; there are 123 occurrences in
src/lib-cuda. Neither is the blast radius, because one entry point can contain
several and because the launcher edge is not counted at all. The unit that
matters is exported entry points whose body is truncated, measured per
object by walking each object's call graph to the exit stub:

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

Every entry point of the public CUDA API. Not "every dispatch" — the whole
surface.

And the compiler had already said so

tests/impl-cuda/compile_probe_mesh.cu has carried this since the mesh port:

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. … it is an artifact of the probe technique, not a defect.

It was not an artifact. That is #150, reported correctly by nvcc and dismissed;
an explicit instantiation was the one thing in the tree that made the check
fire. Corrected in place, and the technique is now used deliberately.

The fix

impl/kernels/utils.h's entire first half — swap, square, sqrt, pow,
min, max, abs, sign, mod, typed_prod ×2, prod ×2, fillfrom ×4,
fill ×2 — was inline FF_CUDEV. All 28 become FF_CUHOSTDEV. Nothing in the
header is device-specific, and its second half (StaticValue and its 60-odd
operators) was already FF_CUHOSTDEV, so the header is now uniform. Device
codegen is unchanged; the host side gains one inline function per used
instantiation.

How the class was found, three ways that agree: (a) nm -u | grep -w exit over
every built object — an undefined exit has exactly one source here, since
nothing in the tree calls ::exit; (b) grep for unqualified calls to each
utils.h name from FF_CUHOST context, which is what turned up the launchers
and build_tree; (c) FF_CUHOST appears 129 times in include/, few enough to
read.

The guards — worth more than the fix

1. tests/impl-cuda/compile_probe_hostdev.cu — calls every helper from a
   plain, non-template__host__ function, i.e. the shape nvcc does check.
   Against the pre-fix header it produces 22 errors; against this branch it
   compiles in 2.5 s. It also explicitly instantiates one launcher per
   affected header, which re-checks the whole launcher body in the device pass.
   Picked up automatically by compile-probe-cuda (PROBESRC is a wildcard).
2. tools/check-cuda-host-stubs.sh — new build-cuda step over
   build/obj/lib-cuda/*.o, fails on any undefined exit. Runs in 2 s and is
   the catch-all for edges the probe does not enumerate.

Nothing that exists today could have gone red: the damage is intra-TU, so
--no-undefined and ldd -r both pass; front-end instantiation is unaffected,
so the FFMEM budget does not move; and the CPU suite cannot see it because
both macros are empty without nvcc.

Cost — measured, and the memory gate does not trip

build-cuda (index32), this branch vs run 32388984092 (main at b0463bc),
same runner image, same flags:

module	peak RSS before	after	Δ	nvcc s before	after
reg_flow	13,604,176 kB	13,605,772 kB	+0.01%	1043.8	1069.3
reg_field	8,523,624	8,630,240	+1.25%	781.4	848.9
reg_field_rls	7,417,244	7,416,756	−0.01%	699.2	745.0
resize	2,100,696	2,100,088	−0.03%	194.9	210.3
reg_flow_rls	1,992,692	1,991,952	−0.04%	163.7	159.5
pushpull_backward	1,580,620	1,580,596	0.00%	162.9	191.8
pushpull	1,555,596	1,553,024	−0.17%	158.0	177.2
restrict	1,361,408	1,361,584	+0.01%	117.8	136.1
distance	843,728	1,041,956	+23.5%	130.1	146.6
posdef	407,556	453,184	+11.2%	44.1	47.6
splinc	439,776	439,176	−0.14%	68.6	68.4
.so link	439,164	541,332	+23.3%	1.5	1.5

Compile step wall: 34m24s → 35m45s (+3.9%). reg_flow stays at 12.98 GiB
against the 14 GiB FF_MEM_BUDGET_KB, with the same ~1.02 GiB of headroom.

Why the peaks barely move, and why that is not evidence the fix does nothing.
FFMEM is ru_maxrss — a maximum over the nvcc process tree. Restoring the
deleted code adds work to the host pass (cc1plus). For the small modules
the host pass is the peak, so it shows: distance +23.5%, posdef +11.2%, the
.so link +23.3%. For the regularisers the peak is cicc/ptxas in the
device pass, which this change does not touch at all, so a real increase in
host work is invisible to that metric — it surfaces instead in compile time
(reg_field +8.6%, pushpull +12.1%, restrict +15.5%, pushpull_backward +17.7%)
and in object size (locally, at identical flags: distance 7.20 → 12.30 MB,
resize 25.14 → 29.79 MB).

The uncomfortable corollary: FF_MEM_BUDGET_KB could never have caught this
defect and cannot confirm its repair. It measures the pass that was always
working. Worth noting next to the #147 discussion — the split is still the right
call, but not because of anything here.

The same across all 11 modules on a local rig (all-Dynamic policy, -O1,
FF_INDEX32=1, so absolute values differ from CI; the ratios are the point):
peaks move ≤0.03% for resize/restrict/splinc/reg_field/reg_field_rls/reg_flow/
reg_flow_rls, and +20% to +52% for distance/posdef/pushpull/pushpull_backward —
exactly the modules whose peak is the host pass.

CPU path — unchanged, proved not assumed

FF_CUDEV and FF_CUHOSTDEV both expand to nothing outside nvcc, so all ten
src/lib-cpu/*.cpp preprocess byte-for-byte identically before vs after
(g++ -E -P, ten of ten). And the gate itself:

tools/test-baseline.sh --tree . --legs default,lib
default leg: 13 suites, 59,886 checks, 0 failures
+ 2 hub suites (14 checks)
15 rows, diffed row-for-row against tools/test-baseline.expected -> identical

(--check insists on --legs all, so the comparison is done by diffing the
default/lib rows, as #147 also did.)

Green

build-cuda (index32), build-cuda (index64), compile-probe-cuda, all six
test-cpu legs, test-hub, codespell. Both new gates ran and passed.
lint (clang-format) is continue-on-error and unreadable (#89) — not chased.

No GPU in CI

Nothing here is a runtime claim. What still needs real hardware: that the
restored dispatch selects the right instantiation for each dtype/dim/bound, that
the launchers launch and the kernels compute correctly, the CUDA stream
plumbing, the atomics, and the mesh sdt path. This PR establishes only that
the code is present and reachable — which it was not, anywhere, for any entry
point.

Adjacent

#154 also edits .github/workflows/ci.yml and CLAUDE.md; the two touch
different regions but will want ordering. #147 is gated on this fix and is
unaffected by it in its own measurements.