cuda(mesh): add the sdt_naive host launcher; zero the vertex normals

[balbasty]

Refs #5. Does not close it — the acceptance bar there is real-hardware
validation, and nothing in this PR has ever been executed (see "What remains
unvalidated" below).

What was verified first

Both gaps named in #5's residual list are still present on main
(include/fastfields/impl/cuda/distance_mesh.h):

• the precomputed-tree/normals sdt overload ends in
  throw std::logic_error("distance_mesh::sdt (precomputed tree) not implemented");
• sdt_naive_kernel exists as a __global__ but has no CUHOST launcher, so
  dt's signed+naive branch threw.

1. sdt_naive host launcher (new)

Structural mirror of cpu-impl's sdt_naive + build_sdt_naive
(include/fastfields/impl/cpu/distance_mesh.h): gather the mesh into
contiguous device buffers (honouring the caller's real input strides), copy to
host, build the face/vertex/edge pseudonormals with the shared kernels
builder, upload, size the grid from prod(size, nbatch) — the batch element
count, not the rank — and launch sdt_naive_kernel on the caller's stream.

No new mathematics: the per-element work is MeshDist::signed_dist_naive from
impl/kernels/distance/mesh.h, the same function the CPU suite covers. Only
loop-and-launcher glue is new. No BVH, no POD tree mirror, no traversal trace —
none of those exist on the brute-force path, which is exactly why this is the
reference sdt should be checked against on real hardware.

sdt_naive_kernel gains nearest_vertex / stride_nearest, which
build_sdt_naive has and it did not; without them the CUDA naive branch could
not fill an output the CPU branch fills. Both are guarded on the pointer, as in
sdt_kernel, because ff::dt_mesh leaves them null by default.

2. Vertex normals were accumulated onto uninitialised memory (fix)

In the existing, already-dispatched sdt. MeshDistUtil::build_normals
accumulates into the vertex normals — normvertices[v].add_(normal) once
per incident face, normalised at the end — which is why cpu-impl allocates them
with new scalar_t[...](). The CUDA launcher used allocHost
(cudaMallocHost), which does not zero. Face normals (copy_) and edge normals
(built in a local map, then copy_) are assigned, so only that one buffer was
affected. Silent failure mode: it perturbs the vertex/edge pseudonormals, i.e.
the sign of the returned distance near vertices and edges — not something
compile+link could ever have caught.

3. The precomputed-tree sdt overload: still a throw, with the reason stated

Not implemented, deliberately. The signature does not carry what a correct body
needs, and inventing the missing contract silently would have produced a
launcher that compiles and cannot be called correctly:

1. const void * tree — sdt_kernel needs a device array of the POD
   DeviceNode mirror, not the polymorphic host MeshDist::Node[] that
   build_tree produces. A void * cannot tell those apart (that ambiguity is
   what fastfields-cuda-impl#44 removed from sdt_kernel by giving the
   parameter a real type), and nothing in the tree hands a caller a
   DeviceNode array — flatten_tree + upload are private to the other sdt.
2. void * treetrace / treesize — the trace is per lane and interleaved
   across lanes, so the buffer must be
   GET_BLOCKS(numel) * CUDA_NUM_THREADS * treesize bytes: a size that depends
   on the launch configuration this function itself picks.
3. faces must be the BVH-sorted face list — build_tree reorders faces in
   place and its leaves index the sorted order.

The TODO(host-launcher) comment (which recorded only the previous bad attempt)
is replaced by the above plus the two ways out: give it a real signature and an
in-repo caller by making sdt delegate to it — the cpu-impl sdt / build_sdt
split — or drop the overload, which has no caller, no producer and no test.
Both are design changes with a reviewable blast radius, so neither is made here.

How CI actually exercises this

tests/impl-cuda/compile_probe_mesh.cu now calls M::sdt_naive directly for
all four (dim × dtype) combinations the cuda-lib dispatcher can select, plus the
null-nearest_vertex shape. probe_dt already covered all four
(_signed × naive) flag combinations, so dt's signed+naive branch was
instantiated before this PR and now resolves to a real launcher rather than a
throw. An uninstantiated template is not compiled, so without that addition a
green build would prove nothing about the new code.

Path filter: this touches impl/cuda/ and tests/impl-cuda/ only, so
build-cuda + compile-probe-cuda run and the CPU legs are correctly skipped.
make test-lib-cpu was run locally to confirm the gate is unmoved — result in a
comment below.

What remains unvalidated

Everything at run time. There is no GPU in CI, so the evidence here is nvcc
accepting and linking the code, on the standing reasoning that the voxelwise
math is shared with the CPU backend and covered by test_distance_mesh.cpp.
Not verified by this PR: the tree walk, the atomics, the stream plumbing, the
normal-buffer lifetimes, and — for the new launcher — whether sdt_naive and
sdt actually agree on a real device. That agreement is #5's acceptance bar and
is still open.

Unrelated and untouched, per #80: src/lib-cuda's MODULES omits
posdef/resize/restrict/splinc. Nothing here changes MODULES or the
link flags.

---

Generated by Claude Code

[balbasty]

CI status

job	conclusion
build-cuda (compile + link)	success
compile-probe-cuda	success
lint (codespell)	success
test-cpu × 4, test-hub, sanitize	skipped (correct — this PR touches impl/cuda/ and tests/impl-cuda/ only)
lint (clang-format, changed lines)	failure — advisory, see below

The CPU gate was run locally instead, since the path filter (correctly) skips it here: make test-lib-cpu CXX=clang++ on this branch gives 59,886 checks across 13 suites, 0 failures — unmoved.

The first build-cuda attempt was cancelled at the 120-minute job timeout, and that was not this diff. It spent 19:11→20:46 UTC (95 min) inside apt-get install nvidia-cuda-toolkit clang — the same slow-mirror episode the test-hub timeout comment in ci.yml describes — leaving ~25 min for a build that needs ~35. Compilation itself was healthy: distance.cpp (the TU that instantiates the new launcher) compiled with warnings only, and the job was killed later, in pushpull_backward.cpp. The re-run took the normal ~35 min and passed.

clang-format — real on my lines, and not fixable without restyling

Not a fluke and not the cancellation: it re-ran and failed on the diff. What it wants is to convert the code I touched from the file's hand-column-aligned style to LLVM style — collapsing the aligned template parameter list, rewrapping the aligned trailing comments, and so on: a ~300-line delta. My added lines deliberately copy the style of the sdt launcher directly above them, including its >80-column aligned parameter comments (which are identical on main).

So this is the tension .clang-format's own header documents ("the tree predates this file … much of the code is hand-column-aligned in ways clang-format cannot reproduce"), which is why the job is continue-on-error: true. I have not restyled: making it green means the new 200-line launcher no longer matches the 1300-line file it lives in. Happy to do it if you'd rather have the green — say the word.

Two things worth knowing about that job while it is in view:

• it can never show its findings. The step is set -euo pipefail with out=$(git-clang-format-18 … --diff), and git-clang-format --diff exits 1 when it would reformat — so set -e kills the step before the echo "$out" in the *) branch ever runs. Every failure renders as a bare Process completed with exit code 1 with no diff. The echo/::error:: branch is dead code. Dropping -e for that one command (out=$(… || true)) would make it useful.
• the sibling failure on docs: dt_mesh returns plain distance, not squared distance #85 was a different thing and is fixed there: two pre-existing over-length @param lines that clang-format reflowed only because an edit landed inside the same comment block. docs: dt_mesh returns plain distance, not squared distance #85 is now fully green.

Audit: are there other launchers with this bug?

The vertex-normal defect here is "a shared kernel accumulates into a buffer the CUDA launcher allocated without zeroing". I swept include/fastfields/impl/cuda/ for the same shape. No second instance. What the sweep covered:

• Every allocation in the layer. Outside the mesh normals, the only data buffers any CUDA launcher owns are distance_euclidean.h's per-lane scratch (v/z/d) and distance_mesh.h's treetrace. The euclidean scratch is write-before-read in the shared kernel and is equally uninitialised on the CPU side (new offset_t[n], no ()), so the two backends agree; treetrace is zeroed by sdt_kernel itself, per lane, before each traversal. Everything else allocated is shape/stride/pole metadata that the copy overwrites in full.
• Accumulate-into outputs.pushpull's push/count/*_backwardout and restrict's out are accumulated into, but they are caller-provided and the pre-zero contract is documented identically on both backends (api/pushpull.h, both restrict.cpps). reg_field/reg_flow accumulate only under the explicit op='+' instantiation. No asymmetry.
• Allocate/copy/free symmetry. Every buffer allocated in a CUDA launcher is freed on both the normal and the exception path; the only two allocations without a matching free are contiguousStrides and copy_faces, which return ownership to their caller by design.
• Signature drift against the CPU reference. Diffed the parameter lists of every same-named function in impl/cpu/ vs impl/cuda/: no parameter exists on one side and not the other (beyond stream). That is the check that would have caught the stride_faces/stride_vertices-ignored bug from fastfields-cuda-impl#42.

One latent fragility found, not a bug today and not touched here: every launcher frees its device buffers immediately after an asynchronous launch, and sdt/sdt_naive do a synchronous, default-stream D2H copy of a buffer written by a kernel enqueued on the caller's stream. Neither is ordered by anything explicit — both are correct only because plain cudaFree/cudaFreeHost synchronise the whole device. If the caller's stream is non-blocking (PyTorch creates its streams with cudaStreamNonBlocking, so legacy-default-stream synchronisation does not apply) and those frees were ever moved, removed, or switched to the stream-ordered cudaFreeAsync, the layer would start racing silently. That is the same "invisible to compile+link" class as the bug this PR fixes; worth its own issue rather than a drive-by.

---

Generated by Claude Code