[teeny] Phase A: distance_{l1,euclidean} dt() takes an anyrank carrier

[balbasty]

Agent:claude-fastfields-to-teeny

Phase A of the tensor-native-boundaries umbrella (fastfields-lib#57) — the impl half.

This is one logical change split across two repos only because GitHub cannot span them: the cpu-lib half is fastfields/fastfields-cpu-lib#74, and neither PR compiles without the other (distance.cpp is dt()'s only caller, and it converts in the same PR-set). Please review them together — same split convention as phase 6's cpu-lib#70 / cuda-lib#39.

Convention reference: fastfields-lib/TEENY-MIGRATION.md §9.

Closes#58

---

What changed

dt() took a decayed (ndim, f, size[], stride[]) tuple and immediately reassembled the tensor its caller had just taken apart. It now takes the carrier itself:

template <classA, typenamescalar_t>
inlinevoiddt(A at, scalar_t w)

Rank, per-line length, per-line stride and the batch offsets are all derived from at (R2/R3) — there is no second source of truth left for anything the tensor already knows. Both distance_l1.h and distance_euclidean.h.

kernels/distance/{l1,euclidean}.h's kernel() is not touched at all (load-bearing for the gate). The euclidean per-worker v/z/d scratch is unchanged (out of scope, per the distance-slice review). The old signatures are deleted, not deprecated.

---

The Phase A gate — all criteria, with evidence

Measured on one machine, clang++ 18.1.3 and g++ 13.3.0, -std=c++17 -DTNY_MAX_RANK=64. A verifiedmake clean ran before each compiler — the script records the post-clean file count each time, and all four reads were post-clean: objects/libs/deps=0 test-binaries=0 (the cpu-lib#56 trap: a green run that silently reused the other compiler's binaries).

1. kernel instantiations byte-identical at -O2 — ✅

Compared as raw section bytes (objdump -s), not merely disassembly. The mangled names are unchanged (the kernel signature did not change), so this is a like-for-like comparison of the same symbol:

symbol	clang++ -O2	g++ -O2	clang++ -O3	g++ -O3
distance_e::kernel<long, float>	byte-identical	byte-identical	byte-identical	inlined ¹
distance_e::kernel<long, double>	byte-identical	byte-identical	byte-identical	inlined ¹

¹ At -O3 g++ inlines the euclidean kernel into the per-line lambda, so it emits no standalone section — and that enclosing symbol (…dt<…>::{lambda(long,long)#1}::operator(), 340/346 instructions) is itself identical before/after, so the kernel's code is still proven unchanged, just as part of a larger identical unit.

distance_l1::kernelnever emits a standalone symbol on either compiler — it is small enough to be fully inlined into the driver. For l1, criterion 1 is therefore subsumed by criterion 2, and the per-line-loop result below is the kernel result.

2. Driver loops: no new instructions in the per-line loop — ✅

The l1 per-line work is a 7-instruction min-sweep body (add / min / inc / store / advance / cmp / jne), emitted once per sweep direction per unrolled copy. Locating every such body by opcode — wherever the compiler chose to put it — and comparing:

build	dtype	base sweep-body sizes	new sweep-body sizes	verdict
clang++ -O2	float	[6,6,6,6,9,9,9,9,9,9,19,19]	same	identical
clang++ -O2	double	[6,6,7,7,8,8,9,9,10,10,19,19]	same	identical
clang++ -O3	float	24 bodies	same	identical
clang++ -O3	double	24 bodies	same	identical
g++ -O2	float	[7,7,7,7]	[7,7,7,7]	identical
g++ -O2	double	[7,7,7,7]	[7,7,7,7]	identical
g++ -O3	float	[7,7,7,7,7,7]	[7,7,7,7,7,7]	same count + same mnemonics; operand-only ²
g++ -O3	double	[7,7,7,7,7,7]	[7,7,7,7,7,7]	same count + same mnemonics; operand-only ²

Instruction count and mnemonic sequence are identical in all eight combinations. Nothing was added anywhere, so there is nothing to explain away or eliminate.

Whole-symbol comparison, matched by structural role (names legitimately change — dt's signature changed):

• clang++, -O2 and -O3: every distance driver symbol is instruction-for-instruction identical — parallel_for<lambda> (serial branch, 253–338 insns), __invoke_impl (195–275), lambda::operator() (184/190), _M_invoke, _M_manager. 18 symbols, all identical, both opt levels.
• g++ -O2: euclidean lambda::operator() (147) identical; l1 parallel branch _M_invoke (89) identical; l1 serial branch — 32-instruction window, mnemonics identical, exactly one operand differs: movss 0xc(%rsp) → movss 0x1c(%rsp), i.e. the spill slot of w moved 16 bytes because the enclosing stack frame no longer builds the carrier locally. Same instruction, different displacement.
• g++ -O3: euclidean lambda::operator() (340/346) identical; l1 parallel branch (105) identical; l1 serial branch same body count and mnemonics, differing only in register allocation and operand order (cmp %r13,%rsi ↔ cmp %rsi,%r13).

² The g++ -O3 serial-branch differences are register-allocation and operand-order choices, not extra work — g++ re-inlines the (now smaller) function differently, and the enclosing function actually shrinks (221 → 184 instructions). The inner sweep is the same 7 instructions.

Method / anti-fooling:objdump -d parsed per symbol; branch and call displacements blanked (they encode placement, not content) while every other operand is compared literally. Negative control: the identical harness run on base -O2 vs base -O3 correctly reports DIFFERS with a diff, so the "identical" verdicts are not an artifact of over-normalisation.

3. Oracle unchanged — ✅

test_distance was 2352 checks / 0 failures at baseline on both compilers (the issues say "2,350"; the true figure is 2352 — 2350 numeric comparisons plus 2 bad-dtype-throw checks). After the refactor, with the original test file, it is still exactly 2352 / 0 on both compilers.

The committed suite then adds 715 deliberate checks (§ below), so the shipped figure is 3067 / 0 on both compilers.

ASan + UBSan (-fsanitize=address,undefined, detect_leaks=1, halt_on_error=1, print_stacktrace=1) on distance.cpp + test_distance.cpp: 3067 checks, 0 failures, no sanitizer reports, exit 0. Built with g++ — this box has no clang libclang_rt.asan, so the clang sanitizer build could not be run here.

4. Whole suite green with unchanged counts — ✅

All 12 test binaries (the issues say 11; there are 12), both compilers, genuine clean between. diff of the full baseline vs post-change check listing yields exactly one differing line on each compiler — test_distance:

binary	baseline	after
test_distance	2352	3067	+715 deliberate (see below)
test_distance_mesh	4622	4622	unchanged
test_distance_spline	704	704	unchanged
test_posdef	5092	5092	unchanged
test_pushpull	326	326	unchanged
test_reg_dispatch	6030	6030	unchanged
test_reg_field	525	525	unchanged
test_reg_flow	1123	1123	unchanged
test_reg_op	186	186	unchanged
test_resize	4130	4130	unchanged
test_restrict	191	191	unchanged
test_splinc	4577	4577	unchanged

Identical table on clang++ and g++; every binary PASSED. distance.cpp also serves the untouched spline/mesh paths, and test_distance_mesh / test_distance_spline are unchanged — those paths did not move.

The +715: what it is, and why it is not drift

test_distance gains test_descriptor_variants, covering the two DLPack descriptor features this path used to normalise by hand:

• 4 cases with strides == NULL (DLPack's compact-row-major shorthand — previously ContiguousStrides), both transforms × both dtypes, each against the brute-force oracle. MIGRATION.md lists strides==NULL as a historical soft spot and no test exercised it at all.
• 4 cases with byte_offset != 0 (previously VOIDPTR), plus a sentinel assertion that the 5 padding elements in front ofbyte_offset are not written — so a mis-folded offset fails loudly instead of reading right and writing left.
• 2 cases combining both.

Total 685 oracle comparisons + 30 pad-sentinel checks = 715, which is exactly the delta.

Crucially, those 715 checks also pass, with the identical 3067 total, against the pre-refactor code — verified by stashing both repos' source changes and rebuilding the new test file against the old ContiguousStrides/VOIDPTR implementation. They pin existing behaviour; they do not encode the new implementation.

Const-correctness (issue item 4 — nothing to do, recorded)

dt mutates in place, so the carrier's element type is non-const here. Phase B's read-only operands will use const-element carriers (§9 R4 records those as probed working on both compilers). The static_assert compares scalar_t against remove_pointer_t<decltype(A::data)>, so a const-element carrier passed here fails with that named message rather than deep inside kernel().

---

Judgment calls (documented — Phase B copies this shape ~40×)

1. Template shape: template <class A, typename scalar_t> dt(A at, scalar_t w).
One parameter per tensor (A, deduced), plus the ordinary deduced parameter per scalar. I deliberately did not re-derive the element type from the carrier and make w a non-deduced element_t<A>:

• It respects R1. dtype is not geometry. R2 says derive geometry from the carrier; dtype dispatch stays in *-lib, so scalar_t should keep arriving from there — exactly as through the old signature. Deriving it here would quietly move a dispatch decision down a layer.
• It scales to Phase B's multi-tensor entries. "One template parameter per tensor, scalars deduce as usual" reads identically whether an entry point takes one carrier or four (posdef's out/inp/hessian); an element_t<A> spelling needs a per-site choice of which carrier defines the scalar.
• Deduction stays total — no explicit template arguments at the call site, so no .template disambiguators on a type-dependent receiver.

The cost is that scalar_t could disagree with the carrier's element type; a two-line static_assert closes it, turning the mistake into one named message instead of a pointer-conversion cascade inside kernel(). Recommended as the Phase B template.

2. offset_t derived, but kernel<offset_t, scalar_t> kept explicit.
using offset_t = decltype(at.size(0)); — the carrier's own offset type, one line, obviously derived. I kept kernel's explicit template arguments rather than letting it deduce from line.extent(0)/line.stride(0): deduction works today, but being explicit guarantees the same instantiation as before regardless of what those accessors return — which is exactly what criterion 1 measures.

3. The size_front<-1>() / parallel_forlong tidy-up the issue asked about.
parallel_for is declared (int64_t, int64_t, int64_t, const F&) and invokes f(begin, end) with int64_t, but the lambdas declared their parameters long. On LP64 those are the same type — hence the byte-identical object code above — but on LLP64 (Windows) long is 32-bit, so the callback silently narrowed a 64-bit line count. Since I was rewriting those exact lines anyway, they now say int64_t: zero codegen change here, a latent correctness fix on Windows. I did not widen scope to other modules' parallel_for lambdas, but Phase B should apply the same spelling wherever it touches one.

---

One finding for the umbrella (a doc bug in §9 R2 — not fixed here)

§9 R2 blesses at.shape(-1) as the spelling for a carrier's trailing extent. That is correct only for a wrapping carrier, and is silently wrong for a copy_meta one — which is what from_dlpack always produces.

as_anyrank(...) without copy_meta builds a _meta_view store whose extent isndim, so shape(-1) wraps to ndim-1 as documented. With copy_meta the store is _meta_store<offset_t, TNY_MAX_RANK> — a fixed-length rank-1 tensor of TNY_MAX_RANK (64 in this build) of which only the leading ndim slots are filled. operator()'s negative wrap is relative to the static extent, so at.shape(-1) reads slot 63 — uninitialised metadata, not the last axis.

distance_euclidean.h needs exactly that quantity, so it reads at.size(at.ndim - 1) with a comment recording why. Phase B hits this immediately (posdef/pushpull read a trailing channel count from carriers built by from_dlpack), so §9 R2's wording should be corrected before Phase B starts. I have not edited the doc from this PR; happy to file it as a Phase-0 follow-up.