Scope the narrowed index arrays, and stop leaking them on the throwing path

[balbasty]

Answers point 2 of the review on #94 — "should the autocast changes be their own
PR?" — with yes. This is the autocast half of that PR, on its own, and it is not
a refactor: it fixes a leak.

What this is worth on its own

A per-call memory leak on the throwing path, reproduced under LSan. That is
true regardless of what happens to the dispatch question, and this PR does not
depend on it.

copy_if_needed / free_if_needed is a manual acquire/release pair, and every
dispatch wrapper in the tree puts code that can throw between the two:

constoffset_t * _size = copy_if_needed<offset_t*>(size, n); // allocates
... as_weights(), newreduce_t[], the impl call ... // can throw
free_if_needed<int64_t*>(_size); // skipped

On the CUDA side the middle section throws by design — every
FF_CUDA_LAUNCH does since #154, and so does every copyToDevice. Measured
with ASan/LSan against the real core/autocast.h, on the reg_field wrapper
shape:

Direct leak of 16 byte(s) in 1 object(s)
Direct leak of 16 byte(s) in 1 object(s)
SUMMARY: AddressSanitizer: 32 byte(s) leaked in 2 allocation(s)

per throwing call on the narrow (int32_t) arm. This is a path user code is
expected to hit — an invalid argument — so it is per call, not once per process.

What IndexArray is

The RAII form of that pair, for the one job it is actually used for: handing a
const int64_t * shape or stride array to a kernel templated on offset_t.

• The throwing path is covered by construction.
• The narrow arm stops allocating. These arrays are nbatch + ndim + 1
  long — single digits in every caller — so the elements live inside the object
  (FF_INDEX_ARRAY_INLINE, 8) and only a genuinely large rank falls back to
  hostNew. That matters most under FF_AUTOCAST_PINNED_HOST, where the
  fallback is cudaMallocHost: a page-locking syscall on the per-call path,
  three to five times per launch.
• The wide arm borrows.IndexArray<int64_t> is specialised to hold the
  caller's pointer verbatim — no copy, no storage — exactly as
  _copy_if_needed<T*, const T*> already does for that case. So under
  FF_INDEX32=0, where both arms name int64_t, the whole class is a pointer
  copy.
• A null source yields a null array and copies nothing, which is what the
  four wgt ? copy_if_needed(…) : nullptr sites in posdef needed.
• Deliberately not convertible to a non-const pointer: nothing in the dispatch
  layer writes through a shape or stride array, and the wide arm aliases the
  caller's memory.

Also: a latent bug in the same header

hostNew / hostDelete throw std::runtime_error on the pinned-host path, and
core/autocast.h never included <stdexcept>. It compiles today only
because every translation unit that reaches it happens to include <stdexcept>
first; a .cu that includes just that header does not:

core/autocast.h(76): error: namespace "std" has no member "runtime_error"

One line. Folded in here because it is the same file and this work is what
found it — I had nvcc available locally this time, which #94's run did not.

Scope

Converts the eight non-regulariser surfaces on both backends — distance,
posdef, pushpull, pushpull_backward, resize, restrict, splinc,
solve_field — 228 call sites across 15 translation units. Every one of
those TUs is fully converted; there are no half-converted files.

Two sites in distance.cpp call .get() explicitly rather than relying on the
implicit conversion, because distance_e::dt / distance_l1::dtdeduce
offset_t from the argument, and a deduced parameter cannot fire a user-defined
conversion. Those are the only two of the 228 that are not a pure declaration
swap.

The regulariser surface is deliberately left out. Its six TUs are being
edited by #158 and by #147's per-(family, ndim) split, and converting them here
would collide with both. copy_if_needed and free_if_needed are untouched and
still used there, so nothing is stranded — the reg conversion is a mechanical
follow-up once those land.

Verification

• clang++ and g++, -std=c++11, -fsanitize=address,undefined — values
  correct on the inline, heap and borrow paths; no leak on the throwing
  path, both compilers.
• nvcc -std=c++14, in the pinned-host configuration the CUDA library builds
  with — compiles.
• normalise-header-guards.py, normalise-include-delimiters.py,
  rename-macros.py, check-cuda-launches.py, codespell — all --check
  clean.
• The full tools/test-baseline.sh --legs default,lib --check run is still in
  progress on a heavily contended machine here; CI's test-cpu is the same
  suite on a dedicated runner and is the check to believe. I will post the
  local rows as a comment when it finishes.

What is not verified: anything about runtime behaviour on a GPU. There is no
GPU here, so whether taking cudaMallocHost off the per-call path changes
upload throughput is unmeasured, and I am not claiming it does. The argument
for that part is that a page-locking syscall per call is worth avoiding on
principle, not a benchmark.

This touches core/, so it correctly triggers the full CI set including both
build-cuda legs. Do not merge until those are green.

---

Generated by Claude Code

[balbasty]

Correcting myself on one sentence in the description, because I checked how
reachable the leak actually is and I overstated it.

I wrote that it is "a path user code is expected to hit — an invalid argument".
That part is wrong. Every FF_CHECK_* throw fires in the exported entry
point, before dispatch reaches the wrapper that calls copy_if_needed:

voidfield_matvec(...)
{
ContiguousStrides _out(out_), _inp(inp_);
FF_CHECK_NO_LANES (out) // <-- throws here …FF_CHECK_SAME_DTYPE(out, inp)
...
FF_DISPATCH_REG(_field_matvec, ...) // <-- … before anything is allocated here
}

So a bad argument does not leak. Nothing is allocated yet. The accurate
statement is narrower:

• CPU — the only throw between the pair is std::bad_alloc out of
  as_weights() / new reduce_t[]. Rare. Effectively unreachable in practice.
• CUDA — the impl call itself throws, and does so for ordinary runtime
  conditions: every FF_CUDA_LAUNCH throws on a bad launch since fix(cuda): route all 39 kernel launches through one checked helper #154, and
  every copyToDevice throws on a failed allocation or copy. Device OOM and
  cudaErrorLaunchOutOfResources are exactly the conditions a long-running
  process hits, and each one now leaks the narrowed shape/stride arrays —
  page-locked ones, under FF_AUTOCAST_PINNED_HOST.

That also answers the obvious reviewer question, "if this is real, why is CI
green on main?" Because it cannot see it:

• the sanitize leg is ASan+UBSan on tests/lib-cpu/, and LeakSanitizer is on
  by default there — but the CPU throwing path needs bad_alloc, which no test
  induces. The validation tests that do assert throws all trip the entry-point
  checks, upstream of any allocation.
• the CUDA path, where the throw is routine, has no runtime coverage at all —
  no GPU in CI, compile-and-link only.

So this sits precisely in the blind spot between the two, which is why it
survived. It does not make the fix less worth having — a per-call leak of pinned
host memory on the device-OOM path is a bad failure mode, and RAII costs nothing
— but "user code is expected to hit this" was too strong, and the honest version
is "the CUDA error paths hit this, and nothing here can observe them".

The 32 bytes in 2 allocations LSan reproduction in the description stands; it
just needed an induced throw to provoke, which is what that harness does.

---

Generated by Claude Code

[balbasty]

Local gate, plus one thing this change does cost that I should put on the
record rather than let a reviewer find.

Gate

$ tools/test-baseline.sh --tree . --legs default,lib
distance default 2352 0 reg_flow default 16347 0
distance_mesh default 4622 0 reg_op default 186 0
distance_spline default 704 0 resize default 630 0
posdef default 4012 0 restrict default 65 0
pushpull default 308 0 solve_field default 452 0
pushpull_backward default 6381 0 splinc default 4577 0
reg_field default 19250 0 lib_device_check lib 5 0
lib_splinc_bound lib 9 0

13 suites, 59,886 checks, 0 failures, plus the 2 hub suites (14 checks),
row for row identical to the default and lib rows of
tools/test-baseline.expected. (--check still exits 1 on a two-leg run
because the recording used --legs all; same artefact as #158, not a failure.)

The cost: object text grows, and that growth is the fix

Instantiation counts do not move — but the emitted code does. clang++ -O1 -fPIC, all-dynamic policy:

TU	instantiations	text before	text after
posdef	502 → 502	218,935	238,523	+8.9%
resize	4,555 → 4,555	2,628,490	2,796,426	+6.4%
splinc	391 → 391	150,122	155,682	+3.7%
solve_field	135 → 135	45,207	46,247	+2.3%

That is not accidental overhead, and it is worth being clear about what it is:
the extra bytes are the unwind cleanup paths that were missing before. The
old copy_if_needed / free_if_needed pair has no destructor, so the compiler
emitted no landing pad — there was nothing to run when an exception passed
through, which is precisely why it leaked. IndexArray has a destructor, so
every call site now gets the cleanup code that frees the array while unwinding.
Paying a few percent of .text for that is the trade this PR is asking for; it
would be suspicious if the size did not move.

Does that reach ptxas?

I do not think so, and it is checkable rather than a matter of opinion. The
growth is in the host code of the dispatch translation units; the
__global__ kernels and everything below them are untouched, and ptxas
processes device code. So the per-module nvcc peak RSS should not move.

The build-cuda job prints exactly that number per module on every run (the
step "Peak nvcc memory per module"), so the two legs on this PR can be compared
directly against the recorded table above MODULES in src/lib-cuda/Makefile:

resize 2.00 GB restrict 1.30 GB splinc 0.42 GB
pushpull 1.48 GB pushpull_backward 1.51 GB posdef 0.37 GB
distance 0.77 GB

Those seven are the modules this PR touches. If any of them has moved
materially, that contradicts the reasoning above and this PR needs another look
before it lands — worth checking when the legs report rather than assuming.
Both are still running.

---

Generated by Claude Code

[balbasty]

Both CUDA legs are green on 50d7cab, which is what this PR was waiting on:

build-cuda (index64) success 22 min
build-cuda (index32) success 41 min
compile-probe-cuda success

Those legs are the real check on this change — they compile all seven converted
CUDA translation units under nvcc, link libfastfields-cuda.so, link the hub
against it, confirm every backend symbol resolves at load time, and run the
exit(1)-stub check from #150. Everything else was green already.

For reference, that timing is normal, not something this PR caused —
index32 runs ~2× index64 on main too (40 and 38 min against 18 and 20, in
runs 32420334642 and 32401730379). Same tree, only FF_INDEX32 differing.

One follow-up commit pushed

7c10367 — whitespace only. Deleting a free_if_needed block left the blank
line that used to separate it from the impl call sitting immediately before the
closing brace:

 posdef::sym_matvec<...>(...);
// <-- this
}

51 of them across the 15 converted TUs. clang-format flagged it and it was a
fair catch, so I fixed that part.

It is provably content-neutral, which matters because the CUDA legs above ran
on its parent:

$ git diff -w --ignore-blank-lines 50d7cab 7c10367 -- src | wc -l
0
$ git show 50d7cab:src/lib-cpu/posdef.cpp | tr -s '[:space:]' ' ' | md5sum
3a34718a473c…
$ git show 7c10367:src/lib-cpu/posdef.cpp | tr -s '[:space:]' ' ' | md5sum
3a34718a473c…

Whitespace-normalised, the files are byte-identical. I held that commit until
the CUDA legs finished rather than pushing it mid-run, since a push cancels
them and they restart from scratch. CI is re-running now and should reproduce.

On the remaining clang-format red

I am not chasing it, and I do not think it should be chased in this PR. Note the
job is already continue-on-error: true, so it does not block. What it wants,
reproduced locally with clang-format-18:

1. Strip the tree's hand column alignment — _size_out (…) → _size_out(…),
   scalar_t * _out = static_cast< scalar_t *> → unaligned. That is
   the majority of the 72 remaining hunks, and it is the thing the job's own
   comment says it will not enforce: "the tree predates it and is
   hand-column-aligned in many places".

2. Reflow a documentation comment destructively. In core/autocast.h it
   would take the code block in the IndexArray docstring and wrap it:

    * const offset_t * _size = copy_if_needed<offset_t*>(size, n); //
   * allocates
   

   which mangles the illustration into nonsense. I am not applying that.

3. # define → #define inside #ifndef — but # define is what the
   surrounding file and core/dispatch.h already use for exactly this
   construct.

So the red is the config disagreeing with the tree's existing conventions, not
new code being untidy. #89 is the tracking issue.

---

Generated by Claude Code

[balbasty]

CI has finished on the tip (7c10367), and it reproduced as predicted — the
whitespace commit changed nothing:

build-cuda (index32) success
build-cuda (index64) success
compile-probe-cuda success
test-hub success
test-cpu (clang-static) success
test-cpu (clang-dynamic) success
test-cpu (clang-cuda-default) success
test-cpu (clang-index64) success
test-cpu (gcc-static) success
test-cpu (sanitize asan+ubsan) success
test-cpu (tsan, grain=1) success
lint (source conventions) success
lint (codespell) success
lint (cuda launch sites) success
detect changed paths success
lint (clang-format, changed lines) failure <- continue-on-error, see above

16 checks, 0 pending, and the only red is the non-blocking clang-format
job discussed in the previous comment.

So both CUDA legs are now green on the branch tip, not just on its parent —
which is what I wanted before saying this is ready, since core/ changes
trigger the full CUDA set and the whole point of those legs is that the CPU
suite cannot see this code path.

Nothing further from me on this one. Ready for review; not merging.

---

Generated by Claude Code