[teeny] Phase A: dt_l1/dt_euclidean import via tny::from_dlpack

[balbasty]

Agent:claude-fastfields-to-teeny

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

This is one logical change split across two repos only because GitHub cannot span them: the impl half is fastfields/fastfields-cpu-impl#60 (which converts dt() to take the carrier), and neither PR compiles without the other. 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#72

---

What changed

dt_l1/dt_euclidean held a complete tensor and disassembled it by hand to feed the impl's raw-args signature. tny::from_dlpack does all three parts by construction — folds byte_offset into the pointer (was VOIDPTR), expands a NULLstrides field to row-major (was ContiguousStrides), and copies the shape/stride metadata into the carrier — so the carrier is built once, in the existing dtype dispatch, and handed straight down:

template <typenamescalar_t = float>
inlinevoid_dt_l1(DLTensor & inp_out, double w)
{
auto at = tny::from_dlpack<scalar_t>(&inp_out);
distance_l1::dt(at, static_cast<scalar_t>(w));
}

ContiguousStridesstays in autocast.h — the spline/mesh paths in this same file still use it (5 remaining uses; Phases D/E, last-user rule), and autocast.h itself is untouched. The l1/euclidean paths already had no int32 arm, so there was no narrowing to delete (R5 effectively pre-applied by the distance-slice review).

---

Preserved behaviour — each item checked, not assumed

1. Exported signatures byte-for-byte unchanged.distance.h untouched; the mangled symbols are identical before/after:

base: T _ZN2ff3cpu12dt_euclideanER8DLTensordi T _ZN2ff3cpu5dt_l1ER8DLTensordi
new: T _ZN2ff3cpu12dt_euclideanER8DLTensordi T _ZN2ff3cpu5dt_l1ER8DLTensordi

2. Error messages and their evaluation order.DISPATCH_DT's body, its float/double arms and its exact rejection string ("only floating point data types are supported") are untouched — the shims keep scalar_tleading their parameter list precisely so the macro's func<float> spelling did not have to change. CHECK_NO_LANES still runs first, on the caller's own descriptor, so the order is unchanged: lanes check → dtype switch throw, and the throw is still reached before any import. from_dlpack's own dtype_matches is a debug _TNY_CHECK that the user-facing throw pre-empts; test_bad_dtype_throws (float16) covers exactly this and still passes.

3. strides == NULL still works — now via from_dlpack's expansion, and now actually tested (it never was; see the +715 below).

4. byte_offset != 0 still works — was VOIDPTR, now from_dlpack, with a direct test plus a sentinel proving nothing before the offset is written.

5. DLPack include order (R7).<teeny/dlpack.h> is included after this repo's vendored "dlpack.h", which was already first in the block — the TU stays on fastfields' DLPack v1.2 throughout (teeny vendors v1.1; same guard, first one wins). I confirmed this was already the natural order and did not reorder it; a comment now records why the order matters.

---

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

---

Bonus: the entry points got materially smaller

Removing ContiguousStrides' conditional heap allocation and its RAII/exception scaffolding shrinks the exported functions (clang -O2, instruction counts):

symbol	base	new
ff::cpu::dt_euclidean(DLTensor&, double, int)	328	172
ff::cpu::dt_l1(DLTensor&, double, int)	318	150

A per-call saving, entirely outside the per-line loop. The metadata copy itself is not new — the old code also built a copy_meta carrier, just one layer down inside dt().

---

Judgment calls

The _dt_l1/_dt_euclidean shims stay, rather than collapsing into the dtype arms. The issue offered either. I kept them because DISPATCH_DT passes the dtype as the leading explicit template argument (func<float>), while the new dt() takes the carrier as its first template parameter — so collapsing the shims would have forced a rewrite of DISPATCH_DT, whose shape and exact rejection message the issue calls behavioural ABI. A shim whose first parameter is scalar_t leaves the macro untouched at zero cost. They are no longer void* re-casts: the import is now their entire body, which is a better justification for their existence than they had before.

A new (benign) failure mode worth stating.from_dlpack debug-asserts that the tensor's device_type matches the carrier's memory space. A kDLCUDA tensor handed to ff::cpu::dt_* now trips that assert in a non-NDEBUG build, where previously it would have dereferenced a device pointer on the host. Strictly an improvement, and unreachable in practice (fastfields-lib dispatches on device before reaching ff::cpu), but it is a behaviour change.

---

Generated by Claude Code