DESIGN.md

Design

How ArrowMetal is put together and what limits it today.

Layers

Python / Rust / Go / C# / R / C++  ──(Arrow C Data Interface)──►  libArrowMetalC (C ABI, handles)
                                                                          │
Swift apps ────────────────────────────────────────────────────►  ArrowMetal (Swift API)
                                                                          │
                                       MetalContext ── pipeline cache ── BufferPool
                                                                          │
                                   MSL kernels generated per element type, runtime compiled
                                                                          │
                                       Metal shared-memory buffers (page aligned, unified memory)

• Buffers are the Arrow columnar layout verbatim: validity bitmap + values (+ offsets for variable-length types). The same bytes serve CPU consumers (C Data Interface) and GPU kernels (C Device Interface): copy-free out always, and copy-free in when the producer's buffers are page aligned — one copy otherwise, which the importer reports (ImportResult.zeroCopy).
• Copy-free import wraps a producer buffer with makeBuffer(bytesNoCopy:) after checking that every page of it is mapped and readable. The check asks the VM map for the regions that cover the buffer (mach_vm_region, one call per region), so its cost does not grow with the page count; see Chunked columns and the import for what it replaced.
• Kernels never assume a length multiple of anything: every buffer is padded to a page so trailing 32-bit bitmap words are readable; kernels bounds-check the last word/element.
• Nulls ride along as bitmaps. Element-wise ops share the input's validity buffer (zero-copy); binary ops AND the two bitmaps on the GPU; compaction repacks.

Where the time goes today (M4 Max, 50M rows)

Source: swift run -c release arrowmetal-bench, recorded 2026-09-06 on the Apple M4 Max (16 cores, 64 GB unified memory, Darwin 25.6.0) that every table on this page was measured on; the dated rounds are in BENCHMARKS.md.

Kernel	Passes over data	Achieved	Bound
sum / min / max	1 read	290-375 GB/s	memory
compare	1 read, 1/64 write	375 GB/s	memory
multiply / cast (vectorised)	1 read, 1 write	375 / 190 GB/s	memory
filter (count, scan, scatter, pack)	2 reads, 1 write	138 GB/s	passes
take (random gather)	random reads	85 GB/s	latency
group-by, privatised (K <= 1024)	1 read + TG atomics	see benchmarks	atomics
grouped min/max (any width)	2 reads + 32-bit atomics	see below	memory
counting sort by group id	2 reads, 1 write	see below	memory

Device peak on M4 Max is ~546 GB/s. Single-pass kernels sit at 55-70% of peak; the rest is command buffer setup and the CPU wait.

Group-by over string keys: a hash table instead of a sort

GroupByKeys turns a key column into dense ids 0 ..< K before any aggregate runs. For utf8 and binary columns that used to mean the GPU string dictionary: hash every row to 64 bits, argsort the hashes, mark run boundaries by comparing the bytes of adjacent sorted strings. Correct, but the argsort is eight radix passes over 50 million 64-bit keys — ~185 ms — and it costs the same whether the column holds a thousand distinct keys or ten million. Real group-by keys are low cardinality, so the sort was paying row-count prices for distinct-count work.

Kernels/StringHashTable.swift replaces it with an open-addressing table sized to the cardinality:

1. Hash every non-null row to 64 bits in one pass over the bytes (sht_hash64).
2. Estimate the distinct count by building a throwaway table over a 1/64 slice of the hash space — a row takes part only when (h >> 40) & 63 == 0. Each distinct string makes that decision once, so the occupied-slot count times 64 estimates the cardinality no matter how skewed the row frequencies are (a row sample would not: it sees the frequent keys and misses the rest). The estimate only sizes the real table, and being wrong costs a retry, never a wrong answer.
3. Build the table with three slots per estimated distinct value: linear probing, slots[s] = row + 1 so every atomic stays 32-bit (Metal has no 64-bit atomic add; it has 64-bit atomic min and max only, see UPSTREAM.md), and the key of an occupied slot is the string of row slots[s] - 1. Equality is decided by comparing the bytes of the candidate against that representative, the same rule is_in uses, so a 64-bit hash collision costs one extra probe and can never merge two different strings. Each row also records the slot it landed in and folds itself into that slot's lowest row index with an atomic_fetch_min guarded by a plain load.
4. Rank the occupied slots with the same GPU scan unique() uses, compact one representative row per slot, and relabel into first-seen order with two argsorts over the distinct count.

Two design choices are worth stating plainly. Nothing is cached beside the slot — not the hash, not a key prefix — because MSL only guarantees memory_order_relaxed: a companion array could be read before its writer had published it, and a probe that then walked past its own bucket would give one string two group ids. Reading hashes[slots[s] - 1] cannot go stale, because the hashes are written by an earlier kernel and never change. And because byte comparison decides equality, this path needs no re-hash retry and no host fallback, unlike the sort path it replaces. The only retry is a table that turned out too small; the last attempt gets two slots per row and an unbounded probe budget, so it cannot fail.

The byte comparison was measured against deciding equality on the 64-bit hash alone: 17.5 vs 13.7 ms at a thousand distinct keys, 22.9 vs 17.0 ms at 100k, and no difference at all at 10 million (where the random slot access dominates). A 20-30% saving is not worth an answer that is only correct to the birthday bound, so there is no unverified mode.

The result is identical to the sort path — same group ids, same group keys, same first-seen dictionary order, same null semantics (a null key still forms its own group, id K) — which is what Tests/ArrowMetalTests/StringHashTableTests.swift asserts, case by case, against both the host hash map and the old path.

Measured (M4 Max, 50M rows, 12-byte utf8 keys, best of 5), 2026-09-06

Group-by sum over a utf8 key column, key mapping included. Metal from Swift; Polars 1.44 (16 threads), pyarrow 25 and the 16-core Swift hash group-by on the same buffers.

distinct keys	Metal before (sort)	Metal after (hash table)	pyarrow	Polars	16-core CPU
1,000	203.3 ms	20.4	58.0	208.1	117.7
100,000	231.5	27.0	194.2	221.4	303.8
10,000,000	322.7	154.1	3387.6	751.7	2336.1

dictionary_encode on the same column, which is the mapping plus the first-seen relabel and the gather of the distinct strings:

distinct keys	sort path (old)	hash table	pyarrow dictionary_encode
1,000	275.8 ms	20.9	422.3
100,000	329.0	26.1	603.0
10,000,000	438.9	268.1	4500.4

At a thousand keys the whole table is 4096 slots — 16 KB — so every probe is a cache hit and the pass is memory-bound on the key bytes. At ten million the table is 32 M slots (128 MB) and the cost becomes the random slot access, which is why the margin narrows from 10x to 2x. The 1000- and 100k-key cases are now 2.8x and 7.2x faster than pyarrow, where the sort path had been 3.5x and 1.2x behind it, and they cost the CPU about 4 ms against pyarrow's 700-2000.

Integer, boolean, temporal and dictionary key columns still take the range path (~22 ms at 50M rows); they never enter the hash table when they are group-by keys.

The same table for the distinct-value functions

unique, value_counts, count_distinct, mode and dictionary_encode over a primitive column had the same problem for the same reason: all five began with an argsort of every row, so all five cost the same whatever the cardinality, and all five were behind the CPU libraries that use a hash table (0.11x to 0.71x of pyarrow/Polars/pandas at 10M rows before this change). Kernels/HashTable.swift is the string table with the byte comparison removed — the key is the 64-bit value, so equality on it is exact — and the estimate, the growth retry, the rank scan and the id pass are literally the same code.

Turning a column into that key is one kernel: integers widen (sign-extended, so the map is injective), and floats are normalised first — every NaN to one bit pattern, -0.0 to 0.0 — so bit equality means Arrow value equality, which is what unique's sort path achieves by normalising before the sort. The order does not change either: the distinct values come back ascending, because the last stage argsorts the K representative values. Sorting K values instead of n rows is the whole trick. count_distinct skips even that — it is the occupied-slot count, so it never gathers or orders anything.

Because the table keeps the lowest row per slot, first-appearance order is available for one argsort of K int32s (HashGroups.firstSlotOrder, exposed as HashDistinct.firstRows) rather than another pass over the rows, which is what order="first_appearance" uses.

The threshold is 1 << 16 rows: below it the sort is a handful of small passes and the table's extra round trips do not pay for themselves. ARROWMETAL_NO_HASH=1 forces the sort path in a release build, which is how the before column below was measured.

M4 Max, 50M int64 rows, Swift, best of 3, 2026-09-06 (the "before" column with ARROWMETAL_NO_HASH=1):

function	distinct	before (sort)	after (hash table)
unique	1,000 / 100k / 10M	142.7 / 147.2 / 303.9 ms	12.3 / 14.5 / 89.7
value_counts	1,000 / 100k / 10M	141.2 / 235.5 / 443.3	14.2 / 17.7 / 110.6
count_distinct	1,000 / 100k / 10M	142.3 / 282.3 / 265.0	9.6 / 10.1 / 42.5
mode	1,000 / 100k / 10M	142.2 / 199.7 / 278.0	13.9 / 17.9 / 121.6
dictionary_encode	1,000 / 100k / 10M	162.3 / 295.5 / 277.2	12.8 / 16.5 / 108.2

Against the CPU libraries on the same buffers, called from Python through Benchmarks/python_gpu_bench.py (M4 Max, 50M rows, wall ms), from Benchmarks/results/python_gpu_bench_50000000_2026-09-07.txt:

function	distinct	ArrowMetal	Polars	pyarrow	pandas	vs the best of them
unique	1,000	15.97	111.03	85.34	75.20	4.7x
	100,000	19.00	123.60	77.98	82.21	4.1x
	10,000,000	100.55	165.92	658.25	620.09	1.7x
value_counts	1,000	17.10	116.33	139.61	140.09	6.8x
	100,000	20.49	320.30	192.25	194.27	9.4x
	10,000,000	109.78	1532.52	1446.67	1496.30	13.2x
count_distinct	1,000	9.59	116.06	81.68	85.50	8.5x
	100,000	9.83	119.60	71.23	76.89	7.2x
	10,000,000	41.29	150.05	714.21	657.65	3.6x
mode	1,000	26.52	58.67	21.49	—	0.8x
	100,000	30.36	70.99	549.52	—	2.3x
	10,000,000	166.28	383.45	988.16	—	2.3x

Two rows need a note. unique at ten million distinct is 1.7x Polars because ArrowMetal orders its output and Polars does not: part of its 100.55 ms is the argsort of the ten million distinct values, which a hash-order unique would not pay. And mode is measured through a Python binding that computes it twice (_mode in python/arrowmetal/__init__.py calls the reduction once for the value and again for the count), so its row carries two reductions. At 1,000 distinct pyarrow's mode is ahead because the values span a narrow range and it counts into a direct-indexed table rather than a hash map.

Strings in two layouts: offsets + bytes, and views

A string column is held in one of two Arrow layouts (MetalStringArray):

• offsets + bytes (utf8 / binary): int32 offsets, one more than the rows, into one data buffer.
• views (utf8_view / binary_view, Sources/ArrowMetal/StringView.swift): 16 bytes per row and any number of data buffers. The first four bytes of a view are the length. A string of 12 bytes or fewer sits in the other twelve; a longer one keeps a 4-byte prefix there, then the index of a data buffer and its byte offset inside it. This is Polars' own layout for String and Binary columns, with data buffers that grow to 16 MB each.

Import. The C Data import takes vu / vz as they are: the views and every data buffer are wrapped in Metal shared buffers without a copy. A view buffer or data buffer does not have to start on a page: MetalArrowBuffer.wrapCovering wraps the pages that hold it and binds the buffer at its offset into the first one (the kernels read it byte-wise or 16 bytes at a time, never in page-sized units), so Polars' buffers, whose starts are not on a 16 KiB page, import without a copy. The validity bitmap keeps the page-aligned rule, because bitmap kernels read it in 32-bit words; it is copied when it is not page aligned or carries a bit offset. The import also totals the bytes of the non-null rows in one GPU pass over the views; a column over 2 GB is refused, the limit the large_utf8 import has. Export hands the same buffers back as vu / vz, so a Polars column crosses in both directions without a copy. An array with more than 64 data buffers (an arrow-rs or pyarrow concatenation of view arrays keeps every input's buffers) has its data buffers copied back to back into merged buffers of under 2 GB each and its views rewritten to point into them, since binding thousands of buffers for every kernel costs more than one parallel copy.

How a kernel reads either layout. Every kernel that reads strings reads them through one accessor, row(i, len), which returns a pointer to row i's first byte and its length (StringLayoutSource): StrOff over offsets + bytes, StrView over views. The kernel body is a template over the accessor, and StringLayoutSource.variants instantiates it once per layout (name and name_v; name_ov, name_vo and name_vv for a kernel over two string columns). A view kernel binds the views where the offsets kernel binds its offsets, and a small table where it binds its data: {0, count} followed by one {GPU address, byte size} per data buffer (MTLBuffer.gpuAddress), so a view that points into any buffer is one indexed load away. The encoder is told about the data buffers with useResources, since nothing binds them directly. The accessor checks every out-of-line view against its buffer's size and reads a view that points outside it as the empty string, so a malformed view, or the unspecified view under a null slot, cannot make a kernel read out of bounds (StringViewTests.testMalformedViewsReadAsEmpty).

What reads views. Lengths, hash32, the four pattern predicates, str_eq against a scalar and against another column, is_in / index_in (either side a view), the gather behind filter / take (the result is offsets + bytes), slice (the result stays a view), the sort keys, the string hash table (group-by keys, dictionary_encode, unique, value_counts), the Unicode and ASCII case transforms, the trims, replace, repeat, slice_codeunits, the pads, str_reverse, count_substring / find_substring, the ascii_is_* / utf8_is_* predicates, utf8_center / utf8_replace_slice, the fused expression kernels (am.query, the engines' filters and string predicates), concatenation of view columns (join keys, unions: the views and data buffer lists are appended and the buffer indices of the later parts moved, no string byte is copied) and every host-side row pass (regex, normalisation).

What converts. Anything else reads offsets / data, and the first such read converts the column on the GPU in a command buffer of its own (lengths, scan, byte copy) and keeps the result, so a column is converted at most once. That is str_concat, the splits, match_like, strptime, the byte-counting pads and utf8_zero_fill, binary_slice / binary_reverse / ascii_reverse, casts from strings and the IPC, Parquet, CSV and JSON writers. MetalArray.string_layout reports "view (converted)" for a converted column, am.string_view_conversions() counts conversions process-wide, and ARROWMETAL_TRACE_VIEW_CONVERSION=1 prints the call stack of each one.

Group-by

Dense group ids 0 ..< K come out of GroupByKeys; everything below aggregates over them. There are four shapes, and the point of the current design is that none of them sorts the key column.

Atomic accumulation (Kernels/GroupBy.swift, GroupBySource.swift) is one linear pass. For K <= 1024 each threadgroup keeps a private table in threadgroup memory and merges it into the device table once; above that the updates go straight to device memory. 64-bit sums are a pair of 32-bit atomic adds with an explicit carry, because Metal has no 64-bit atomic add (only min and max, UPSTREAM.md).

Correctly rounded Float64 sums (Kernels/GroupSumExact.swift, GroupSumExactSource.swift) give sum and mean over Float64 values without any group order. A correctly rounded sum does not depend on the order of its terms, so the rows are added where they lie, in two passes of 32-bit atomics. Pass one takes each group's largest exponent (an atomic max of the 11-bit field), counts its values and notes NaN, ±inf and +0.0 in a flags word. Pass two writes each value as an integer in a fixed-point window whose unit is 2^-G of the group's largest exponent's last bit and adds it to a 128-bit two's-complement accumulator, one 32-bit word at a time with the carry (or borrow) taken from each word's own read-modify-write, so the words hold the exact sum modulo 2^128 however the rows interleave. G = 74 - bitlength(rows): 53 significand bits, G guard bits and bitlength(rows) of headroom stay below the sign bit, so no sum can overflow it. A value whose last bits fall below the window is truncated and counted. The finalize, a thread per group, rounds once to nearest-even: with nothing truncated the accumulator is the exact sum, and the mean is the exact sum divided by the count in a long division, rounded once. With d values truncated the exact sum lies strictly between W - d and W + d units; the result is taken when both ends round to the same double, and the rare group whose rounding the window cannot settle (or that holds a NaN other than the canonical one) is summed exactly on the host, from its rows alone. Both come out of one call, which the GroupBy keeps, so a sum and a mean of the same column cost one pair of passes. Up to 1,024 groups each threadgroup accumulates in threadgroup memory and merges once, as the atomic accumulation does.

Sort-free extremes (Kernels/GroupByExtrema.swift) is how min, max and hash_min_max avoid the missing 64-bit atomic. Every element type maps order-preservingly into a 64-bit unsigned key. Pass one takes the extremes of that key's high word with 32-bit atomics; pass two takes the extremes of the low word among only the rows whose high word already equals the winner. Some row attains the winning high word, and among those the smallest low word is the overall minimum, so the answer is exact. Types four bytes wide or narrower skip pass one — their whole key is the low word. A group's four words sit together in one 16-byte entry, a row reads them before it writes and issues an atomic only when it improves on what it read, and no count is kept: a group has a value exactly when its minimum key is at or below its maximum key, which the empty identities never are. Both extremes come out of one call, which the GroupBy keeps, so a min and a max of the same column cost one pair of passes. first / last ride the same kernel over a masked row index.

Counting sort by group id (Kernels/GroupOrder.swift, GroupOrderSource.swift) replaces the argsort for the aggregates that need each group's rows together — list, distinct, product, the central moments, the order statistics. Histogram, scan, scatter: one pass over the keys rather than four over keys and payload. Two scatters share the scan:

• chunked, when a per-block histogram fits the budget. A "block" is one simdgroup, so a row is ranked against 32 lanes rather than 256, and each block owns a reserved slice of every group's run and fills it in row order. Stability is exact and needs no atomics.
• atomic, when the group count is large. One atomic bump per row places the rows, then cs_fix (a thread per group) or cs_fix_tg (a threadgroup, bitonic) sorts each run back into row order. Cheap exactly while runs are short, which is the case that forced this path; a longer run falls through to the argsort, so no input is worse off than before.

The ordering is cached on the GroupBy, so N aggregates over one grouping pay for it once.

Segmented reductions come in two shapes. _wide gives a threadgroup to a group and tree-reduces in threadgroup memory; _narrow gives a single thread to a group. Ten million groups of five rows want the second one — the first would launch a quarter of a billion threads to do fifty million adds.

Grouped variance, stddev, skew, kurtosis (Kernels/GroupMoments.swift) run the two-pass shifted algorithm over that ordering in true binary64: an exact d_add sum for the per-group mean, then the sums of d = x - mean and d^2 (and d^3, d^4) about it, finished as m2 = (sum d^2 - (sum d)^2 / n) / (n - ddof). Shifting first is what makes it accurate — sum x^2 - (sum x)^2 / n cancels catastrophically on data far from zero — and the (sum d)^2 / n term removes the error in the mean itself. On 200k float64 values offset to 1e9 in 997 groups, the largest relative error against an exact rational reference is 1.9e-16; pyarrow's own answer on the same input is 2.5e-10 off. The Float32 deviations this replaced cost about 1e-5.

Finalizers write the output column and its validity on the GPU. A Swift loop over the group count is invisible at a thousand groups and is the whole cost at ten million.

Before and after (Apple M4 Max, best of five, 10M rows, one group_by per call), 2026-09-06

grouped aggregate	groups	before	after	change
min (int64 values)	1,000	18.7 ms	2.7 ms	6.9x
min	100,000	35.8 ms	3.5 ms	10.2x
min	10,000,000	1189 ms	17.0 ms	70x
max	1,000	18.8 ms	2.7 ms	7.0x
max	100,000	36.0 ms	3.5 ms	10.2x
max	10,000,000	1228 ms	17.8 ms	69x
mean	10,000,000	338 ms	30.3 ms	11x
variance (float64, ddof=1)	1,000	512 ms	6.5 ms	79x
variance	100,000	500 ms	11.1 ms	45x
variance	10,000,000	823 ms	34.8 ms	24x
first	100,000	10.7 ms	4.4 ms	2.4x
first	10,000,000	480 ms	18.1 ms	27x
list	1,000	20.3 ms	6.7 ms	3.0x
list	100,000	20.8 ms	11.6 ms	1.8x
sum	any	unchanged (already the atomic path)
count	any	unchanged
count_distinct	any	GPU hash set over (group, value) (TO_IMPROVE.md)

At 50M rows, against the fastest idiom of Polars / pyarrow / pandas (Benchmarks/results/full_matrix_2026-09-07-parallel.csv): min 3.32x / 4.93x / 4.57x at 1k / 100k / 10M groups, where it was 0.25x / 0.56x / 0.21x; variance 0.96x / 2.52x / 4.06x, where it was 0.05x at 100k; first 7.95x / 5.81x / 5.12x; list 10.85x / 9.72x / 5.70x. Two caveats on those ratios. At a thousand groups every numeric grouped aggregate lands near 3.3x to 4.1x, sum and count included, because a fresh group_by([...]) spends about 7 ms of the 10 ms rebuilding the dense key mapping — reuse the object and the aggregate itself is 3 ms. And count_distinct by key was the remaining shortfall and is now 3.4x to 7.2x the fastest CPU idiom (TO_IMPROVE.md) — the sort of the key column is gone from every aggregate.

Latency (small inputs)

Measured floor on M4 Max: an empty kernel with encode + commit + wait costs about 60-70 µs measured (RESIDENT.md), 110-160 µs per call for sum and filter at 1,000 rows in the matrix's latency family; ten kernels in one command buffer cost ~100 µs in total. So the fixed cost is the round trip, not the kernel, and the only lever is fewer round trips.

Unbatched, every public call is one command buffer. filter already fuses its three kernels into one command buffer with a GPU scan.

Batched execution (MetalContext.batch { }, am_batch_begin/end in C, with am.batch(): in Python): kernels append to one serial compute encoder; results are created immediately with their buffers, but a result whose length the GPU decides (filter) or whose null count needs a bitmap read is marked pending. Any CPU-side read (a reduction, length, nullCount, subscripts, export) triggers a sync point: the open command buffer is committed and waited once, deferred fix-ups run (lengths, null counts, take bounds errors), and a fresh batch is opened so the caller keeps batching. While any batch is open the pool parks returned buffers instead of recycling them, so pending GPU work can never observe a reused buffer. Per-thread batches; nested batch calls join the outer one.

Lengths flow on the GPU. Every kernel takes its element count as device const uint* nPtr. A pending filter result binds the GPU-written total as that pointer and is dispatched at its maximum possible size, so the next kernel (compare, arithmetic, cast, bitmap ops, another filter, a reduction's partial pass) never needs the CPU to know the length. A reduction still syncs once to read its partials.

Float64 without hardware doubles. Kernels/DoubleMath.swift carries a software IEEE-754 binary64 (d_add, d_sub, d_mul, d_div, d_sqrt) over ulong bit patterns, correctly rounded. Sum accumulates with d_add; arithmetic kernels run one element per thread. d_div is a Newton reciprocal seeded by one hardware float division, with the exact 128-bit remainder N - q·D settling the last bit; d_sqrt refines a hardware rsqrt seed with Newton steps in fixed point and settles the last bit with the exact remainder. Both are correctly rounded rather than close, and DoubleMathTests holds them to Swift's own Double bit for bit. add and multiply run at this ~390 GB/s these single-pass rows reach at 50M rows and divide within 15% of it (331 GB/s), so the software arithmetic is all but invisible in a bandwidth-bound query. cast is the one float64 path that still runs on the host, and modulo is the one binary operator with no float64 kernel — both say so in COVERAGE.md.

The transcendentals are a different story, and worth being explicit about. Kernels/DoubleTranscendental.swift (expm1, log1p, logb, hypot, the ten RoundModes) and Kernels/DoublePower.swift (exp, ln, log2, log10, pow) evaluate in binary64 from end to end, over that same software arithmetic — not by narrowing the column to float, calling Metal's own library and widening back, which is what these kernels used to do for about seven correct significant decimal digits out of sixteen. pow is the one that sets the bar: a 1-ulp result needs the product y·log2 x, which reaches 1024 in magnitude, accurate to 2⁻⁶¹ absolutely — more than a double holds — so log2 x is carried as an unevaluated hi/lo pair whose high part keeps 21 significant bits, and y is split the same way so that y₁·t₁ is an exact double. That is fdlibm's layout, and its Remez coefficients and hi/lo constants are reused verbatim; the three logarithms fall out of the same reduction, more accurately than a direct series would give them.

Accuracy is measured, not derived. DoubleTranscendentalTests compares each function with Foundation over 10⁶ random inputs drawn across its whole domain and prints the ulp histogram; these are those numbers, measured 2026-09-06. The measured column is what the run reports; the test asserts a 2-ulp bound on everything but sqrt, which it holds to bit equality, so a regression of one ulp shows up as a changed number here before it fails the suite. The trigonometric family is a separate kernel and a separate budget: measured 2-5 ulp against a 6-ulp assertion (TrigTests, COVERAGE.md).

function	domain sampled	max ulp
sqrt	10⁶ random positive bit patterns drawn uniformly over the exponent range as well as the significand — subnormals, 1e-320 and 1e308 included — plus the perfect squares, their neighbours, every binade and the subnormal range in testSqrtIsCorrectlyRoundedOnAdversarialInputs	0 (bit-identical)
exp	-745.2 to 709.78, plus the subnormal-result and near-overflow edges	1
ln	5e-324 to 1.8e308, plus near 1 and the subnormals	1
log2	same	1
log10	same	1
power	10⁶ random pairs, 5·10⁵ negative bases with integer exponents	1

The C99 edge table — x^0, 0^y, 1^y, (-1)^int, infinity and NaN propagation — matches libm bit for bit, and the _checked twins are the unchecked kernel plus a read-only check pass, so they inherit every value and raise at exactly the same boundaries.

The price is throughput, and it is the cost of the accuracy. At 50M rows on an M4 Max, the old float-detour ln ran at 250 GB/s because it was memory bound; the binary64 one runs at 10 GB/s because it is compute bound on forty-odd software operations per element. That is a 25x drop in throughput for nine more correct digits, and it is the right trade for a library whose whole claim is that a Float64 column means Float64. The float32 kernels are untouched and still take the hardware path.

Async

MetalContext.batchAsync is batch { } without the wait. It records body's kernels into one command buffer exactly as batch does, then commits and hangs an MTLCommandBuffer.addCompletedHandler off it instead of spinning. The deferred fix-ups — afterFlush closures (pending lengths, null counts, take bounds errors), pool.releaseParked(), the openBatches decrement — run on the GPU's completion thread, so the calling thread is released the moment the work is committed. Two forms:

let kept = try await ctx.batchAsync { try amount.filter(where: .gt, 100) }   // suspends, no thread held
ctx.batchAsync({ try amount.filter(where: .gt, 100) }) { result in ... }     // returns immediately

flush is split into detachBatch() (end encoding, unhook from the thread) and finishBatch(_:) (the post-completion fix-ups, which touch no thread-local state); the synchronous path calls both back to back around its wait, the async path puts the completion handler in between. Batches stay per-thread: body runs synchronously, before the first suspension, so the batch is opened, filled and detached without a thread hop. Nested calls join the enclosing batch, as with batch.

Sync points still exist inside body. A reduction, length or nullCount of a pending filter result, a subscript or an export commits the open batch and blocks right there, exactly as in batch { }. So body should return something that does not force a sync — a pending array from filter, compare, take, cast or arithmetic. Its length, null count and contents are resolved by the time the await returns. For a scalar there is an async accessor: MetalArray.sumAsync() (and meanAsync()) records the reduction kernel in an async batch and combines the per-threadgroup partials on the completion path, so nothing blocks. Deferred errors are thrown from the await (or delivered as .failure), not from the recording.

Not available: async accessors for min/max/group-by, and cancellation (a committed command buffer runs to completion).

CPU/GPU router

Every public call without a batch pays one command-buffer round trip (above), and on small inputs that floor is the whole cost. The router is the answer the engine gives to that: for seven operations it has a second implementation, a tight single-threaded CPU loop, and it runs that loop when its crossover table says the CPU is faster for an input of this size (the table and what its CPU side measured are below). The output is the same Arrow array either way, byte for byte.

What is routed. sum, min, max (and mean, which is sum over the valid count), compare (scalar and array), add / subtract / multiply (scalar and array), filter by a mask and the fused filter(where:), and GroupBy.sum over integer values (and GroupBy.sumUnsigned, uint64 kept unsigned, which the C ABI and Python use for uint64) with at most Router.groupBySumMaxKeys (1,024) keys, which is also what hash_sum reaches through the dense-key mapping. Every other operation is unchanged and always runs its GPU kernel.

Where the decision is made. Once, at the top of each routed operation's public entry point (Kernels/Reductions.swift, Compare.swift, Arithmetic.swift, Filter.swift, GroupBy.swift), before anything is encoded: if Router.decide(.sum, self).isCPU { return RouterCPU.sum(self) }. Not in MetalContext.run, which sees an encode closure and neither the operation nor the row count. The code is in Sources/ArrowMetal/Router/.

The decision, in this order (Router.decide):

1. No CPU path, no choice. If the operation has no CPU loop for this input the GPU runs it, as before the router, and the decision records why ("no alternative: ..."). That is Float32 min / max, Float32 arithmetic and the fused Float32 filter(where:) predicate, which compute in hardware float on the GPU, where subnormals are flushed (EVALUATION.md, finding 1), so a CPU loop would not give the same bits; divide, which is not routed; and a group-by sum over more than 1,024 keys.
2. Pending input. An input whose length the open batch is still deciding stays on the GPU.
3. Mode gpu runs the GPU.
4. Inside a batch, the GPU. MetalContext.batch shares one command buffer's floor across every kernel in it, and queued kernels may be producing this very input; a CPU read would have to commit and wait for them, which is what batching exists to avoid. So any open batch keeps the routed operations on the GPU, including under mode cpu.
5. Mode cpu runs the CPU loop.
6. Mode auto (the default) looks the operation up in the crossover table: below the crossover the CPU loop runs, at or above it the GPU. The table was measured on int64 columns; integer columns of every width use those rows, and floating-point columns, which have no measured crossover, stay on the GPU. The arithmetic row was measured on add, and subtract uses it. multiply, whose 64-bit integer form costs more per element on the CPU than an add, has its own row (Router.crossoverRows(arithmetic: .mul)), fitted from the shipped multiply loop (below).

The router does not adjust for residency: every routed call starts from Metal-shared buffers (an imported column is mapped before the operation is called), which is the state the table was measured in. It keeps no state across calls beyond the table in force, the modes and the last decision, and it never times or estimates anything at call time: Router.route, the rule every routed call goes through, is a pure function of the operation, the value type, the row count, the mode, the batch state and the table row.

The table is data generated from the measurement, not typed: Benchmarks/router_table.py --from-check Benchmarks/results/router_check_2026-09-24.csv fits it from a Benchmarks/router_check.py run, which times every routed operation with the router pinned to gpu and to cpu, so the CPU side is the shipped RouterCPU loops, and writes Sources/ArrowMetal/Router/RouterTable.swift, a Swift literal, so there is no resource bundle to load. The step alone would send every size between the last measured size where the CPU was ahead and the first where the GPU was ahead to the CPU, so the script fits the crossover once, offline: inside that bracket both paths are taken as straight lines between the two measured points and the crossover is where they meet. It never leaves the measured bracket. multiply has its own row, fitted the same way from the multiply(int64, 3) rows, because a 64-bit integer multiply costs the CPU more per element than an add. router_table.py --check fails when the committed table no longer matches the results file it names, and python/tests/test_router.py runs it. The values are in the generated file, in Router.crossoverRows(_:) and Router.crossoverRows(arithmetic:) from Swift, am_router_crossover and am_router_multiply_crossover from C and am.router_crossovers() from Python.

Benchmarks/results/router_check_2026-09-24_after_refit.csv is the same check run against the refitted table: auto took the faster path in 71 of 72 cases. The one miss is compare(int64 > int64) at 3M rows, by a margin of 1.10, because the compare row is fitted on the scalar compare. Both runs compare the GPU with a single-core CPU loop; a background process holding one of the sixteen cores does not enter either side.

The table can also be generated from Benchmarks/results/router_2026-09-17.json, the size sweep behind CROSSOVER.md (router_table.py --json), whose CPU side is the bench's own single-core loops in Sources/ArrowMetalBench/main.swift rather than RouterCPU; the generated file's header says which source it came from, and --check regenerates from that source.

A table per machine. python -m arrowmetal.router calibrate [--quick] runs the same sweep on the machine it is run on, fits it with the same code (python/arrowmetal/_router_fit.py, which router_table.py loads too) and writes ~/.arrowmetal/router/<chip>.json. At first use the router takes ARROWMETAL_ROUTER_TABLE (a path, or shipped), else that file for this chip, else the shipped table; python -m arrowmetal.router explain <op> <rows> prints a decision with the table row it came from. CROSSOVER.md describes the pipeline and the JSON format.

The CPU side (Router/RouterCPU.swift) is one generic loop per shape, over the Arrow layout: reduce (a branch-free fold where a null slot feeds the operation's identity), map (arithmetic, every slot including those under nulls, as the GPU computes them), map-to-bitmap (compare, 32 bits per word, bits past the length zero), compact (filter, by 64-bit selection words), and a dictionary loop for the group-by sum (a keyCount-slot table of wrapping 64-bit sums and counts). They are single-threaded; CROSSOVER.md also records the bench's all-core loop for reference. CPUReference stays the tests' oracle; its closure-per-element walk is one to two orders slower than these loops (CROSSOVER.md, cpu-ref).

Byte-identical output. For every routed operation the CPU loop returns what the GPU kernel returns: the same values, the same validity buffer (the scalar forms share the input's bitmap exactly as the GPU does; the array forms AND the two bitmaps), the same null count and the same buffer sizes. Integer arithmetic wraps as the GPU does. Float64 arithmetic is the hardware add, subtract or multiply, which is correctly rounded like the software d_add / d_sub / d_mul, with the software kernels' NaN rule applied when the result is a NaN (the first NaN operand, quieted; otherwise the positive quiet NaN). Float64 min / max compare the GPU's order-preserving keys, so NaN is skipped and -0 and +0 are one key. Float sums reproduce the GPU's summation order exactly rather than adding in row order: thread t of the groups x 256 grid adds its 4-element blocks t, t + grid, ... and then tail element t, each threadgroup folds its 256 accumulators as the kernel's tree does, and the host adds the group partials in order, as finaliseSum does. The sum is therefore the GPU's to the bit, NaN payloads included, with no Accelerate and no tolerance. RouterTests (Swift) and test_router.py (Python) hold both paths to this on every type, on sizes across the word and threadgroup boundaries, with 0%, 10% and 90% nulls, on slices at an offset, with signed zeros, infinities, quiet and signaling NaNs with payloads and subnormals, and on float sums past the point where every GPU thread takes more than one block.

Overrides and visibility.

	Process	One call or one scope	Last decision
Environment	ARROWMETAL_ROUTER=auto|gpu|cpu, read at load
Swift	Router.mode	Router.withMode(.cpu) { try col.sum() } (per thread)	Router.lastDecision (per thread)
C	am_router_set_mode, am_router_get_mode	am_router_set_thread_mode (-1 clears)	am_router_last, am_router_last_reason
Python	am.set_router("cpu"), am.get_router()	with am.router("cpu"):	am.last_route() -> RouteDecision(op, path, reason, rows)

The per-thread override wins over the process mode. Bindings other than Python reach the router through the C functions above; no operation signature changed.

What the router does not do. It does not remove the floor, and it does not touch the operations whose cost is not the floor (memory-bound single passes, software binary64 transcendentals, host regex, views, temporal extraction: TO_IMPROVE.md). It chooses between ArrowMetal's own two paths, never a CPU library. Strings are not routed: the string rows are behind a 12-core Acero, and a single-threaded loop would not change that.

Top-k and order statistics: GPU radix select

top_k, quantile and approximate_median all ask the same question — which rows sit at which rank — and all of them used to answer it by sorting the whole column. Sorting 50M Int64 is eight LSD radix passes and a gather, about 130 ms; the answer needs a few hundred bytes of it. Radix select replaces the sort with two passes over the column, whatever k is.

The algorithm

Everything runs on the order-preserving key: the same map argsort uses (TopKSource.keyMap), so nulls, NaN, -0.0 and the ascending/descending flag behave identically and the answer is exactly argsort(descending:)[..<k], index for index.

1. Histogram (rs_histogram, Kernels/RadixSelectSource.swift). One pass counts the top 8 bits of every non-null row's key. The counts are kept per simdgroup-sized sub-block — one 256-bin private histogram per simdgroup in threadgroup memory, so the only atomic contention is between 32 lanes — and written digit-major as counts[digit * subBlocks + sub]. Summed over sub-blocks that is the global histogram; left per sub-block it is exactly the offset table step 3 needs. That double duty is what keeps the fast path at two passes instead of three.
2. Pick the bin. The host walks the 256 totals to the first digit whose running total reaches k. Rows with a smaller digit are guaranteed winners — there are fewer than k of them — and the rest of the answer lives in that one bin, about n/256 rows for well-spread keys.
3. Compact (rs_scatter). A second pass writes the winners and then the bin's rows, both in ascending row order, into one array of (key, row) pairs. Every winner's key is smaller than every bin key, so a stable sort by key alone turns that array into (key, row) order — the total order — with no second sort key and no 96-bit comparisons. Because each sub-block owns a contiguous slice of the output, the scatter needs no threadgroup barriers: a simd_prefix_exclusive_sum per 32 rows is the whole ranking.
4. Refine. The bin is still ~n/256 rows, far more than the answer, and it would dominate the final sort. So steps 1-3 run again on the compacted array and the next digit — microseconds of work that shrinks the bin by another factor of 256. One refinement round is enough; the loop stops when the bin is under TopK.refineCap (4096) rows.
5. Order. Up to TopK.bitonicCap (2048) survivors are sorted by a bitonic network in a single threadgroup (rs_sort_small, one dispatch, the slot index as the tie-break), which covers every k up to about 1024. Above that the stable LSD radix sort orders them and the first k rows are the answer.

Ties. Heavily tied data can put nearly every row in the chosen bin, which would make step 3 copy the whole column. Step 2 therefore repeats on the next digit without compacting while the bin is over the candidate budget (1M rows). When the key runs out of digits every remaining row is an exact tie, and ties are settled by row order, so the scatter writes only the first k - winners of them: the output is exactly k pairs however many rows share the key.

The same in-place narrowing covers the other skewed case: a float64 key's top byte is the sign plus seven exponent bits, so a column of uniform doubles concentrates in a handful of bins rather than spreading over 256, and a rank in the middle of one of them costs one extra pass before the bin fits the budget.

Nulls. The histogram and the scatter both skip invalid rows, so the selection is over the non-null rows only. When fewer than k rows are non-null the answer has to reach into the null rows, which only the sort places — that case, and k above half the non-null rows (where the compaction would copy most of the column anyway), fall back to argsort.

Order statistics: selecting a value, not rows

A quantile wants the value at a rank, not the rows, so Kernels/RadixSelectValue.swift runs the same narrowing with the winners counted but never written (wantLt = 0). Once the bin is under 4096 keys it is read back and finished on the host, which ends the round trips immediately instead of after another six digits. kthElement(k, largest:) is that search plus the inverse key map; quantile(q) asks for the one or two adjacent ranks that bracket the interpolated position, and because they are adjacent one search usually settles both. approximateMedian() is quantile(0.5).

The reconstructed value is exact for every integer type. For floats it differs from indexing a sorted copy on exactly the two values the key deliberately merges: -0.0 comes back as +0.0 and a NaN loses its payload. Both are values Arrow's total order calls equal, and quantile returns a Double, so its answers are unchanged.

tdigest still sorts: its centroids are defined by the whole sorted sequence, not by one rank, so selection cannot produce the same digest.

Which path runs

condition	path
k > 1024, or n >= 2^19	radix select
k <= 1024 and n < 2^19	per-threadgroup selection (Kernels/TopK.swift) — one dispatch, no mid-flight readback, so it wins where latency rather than bandwidth decides
fewer than k non-null rows, or k > (non-null)/2, or a type with no key map	full argsort

The crossover at 2^19 rows is where radix select's two passes plus one host readback start beating a single selection dispatch. Below it the per-threadgroup kernel is one command buffer; above it, it is bandwidth that matters and radix select reads the column at close to peak.

Numbers (M4 Max, 50M rows)

Called from Python on the same in-process data (PYTHONPATH=python python Benchmarks/python_gpu_bench.py 50000000 5), against pyarrow 25, Polars 1.44 (16 threads) and numpy 2.5. Wall ms, from Benchmarks/results/python_gpu_bench_50000000_2026-09-07.txt. Before radix select, top_k(10 000), top_k(100 000) and both quantiles ran a full argsort or sort of the column.

Operation (50M rows)	ArrowMetal	pyarrow	Polars	numpy
top_k(100), Int64	2.65	24.59 select_k_unstable	67.35	183.58 argpartition
top_k(10 000), Int64	3.02	34.78	62.06	182.37
top_k(100 000), Int64	3.16	154.78	61.29	183.50
quantile(0.5), Int64	2.56	275.87	70.09 median	280.94 median
quantile(0.5), Float64	2.27	353.04	126.01 median	362.31 median

That is 9.3x / 11.5x / 49.0x pyarrow's select_k_unstable at the three k, and 108x / 156x pyarrow's quantile. The target was 3x. top_k is flat in k, because k only changes how many of the compacted candidates the final ordering has to sort, not how much of the column is read — while pyarrow's heap-based select_k_unstable goes from 24.6 ms to 154.8 ms over the same range.

The column is read twice, so the useful bandwidth figure is 2 x 400 MB over the wall time: 302 / 265 / 253 GB/s for the three top-k, against a device peak of ~546 GB/s. The gap is the mid-flight readback — the host reads 256 counts to pick the bin, so the two passes cannot share a command buffer — plus the final ordering, which grows with k.

The CPU stays free throughout: 1.4-1.8 CPU-ms for the top-k calls and 0.9-1.2 for the quantiles, against 24.6-362.3 CPU-ms for the CPU libraries on the same questions (pyarrow, Polars and numpy; Benchmarks/results/python_gpu_bench_50000000_2026-09-07.txt, rows 259-278).

The radix sort's block size is also now adaptive below ~256k rows: a fixed 4096 elements per block left a 20k-element sort — the size top-k's final ordering lands on — running on five threadgroups. Inputs above ~256k rows are unaffected, so the 50M argsort is unchanged.

Index arrays: UInt32 row numbers

Every call that returns row numbers into its input returns them as UInt32 (Arrow uint32, format I): argsort (numeric, boolean, temporal, utf8 and binary), topK, partitionNthIndices, the lexsort, the hash join's index pairs (hashJoin, JoinExtra.indices), the ranks (rowNumber, rank, denseRank, maxRank, rank(tiebreaker:), and the engine's row_number / rank / dense_rank window columns), GroupByKeys.representativeRows() and GroupByKeys.rowIndices(_:_:). UInt32 is the width the Int32 arrays had, so the kernels, the bytes they move and the values they write are the same; it is the type arrow-rs sort_to_indices and Polars' IdxSize use. pyarrow returns uint64 for the same functions, with the same values. Before, a row number of 2^31 or more came back negative in an Int32 array.

Arrays hold at most 2^32 - 1 elements, so every row number fits. A call whose row numbers would pass 2^32 - 1 is refused rather than wrapped, with the error "<op>: row N does not fit the UInt32 index type (row numbers go up to 4294967295)" (Dispatch.checkIndexRows). take accepts Int32, Int64 and UInt32 indices, as Arrow's does; an Int64 index the 32-bit gathers of nested, fixed-size-binary and interval columns would have to narrow is checked against the range first, so it is an error and never another row.

The integer outputs that are not row numbers into the input keep their types: dictionary codes and index_in positions are int32 (Arrow's defaults), group ids are int32 (a group-by holds at most 2^30 groups), inverse_permutation is Arrow's signed int32 and refuses an index column of more than 2^31 rows, list_parent_indices() is int32 and refuses a list column of more than 2^31 rows (listParentIndices64() is int64), and indices_nonzero is uint64, as in Arrow. The hash join takes at most 2^31 - 1 rows per side and returns at most 2^31 - 1 pairs.

Timed against the Int32 build, alternating in one process (Benchmarks/index_type_bench.py, Benchmarks/results/index_type_2026-10-01.csv): an untimed warm-up of at least 100 ms per call, 8 rounds of up to 30 calls, and the three calls whose median moved by more than 1% timed again alone, 20 rounds of 100 calls (marked *). No call is more than 1% slower in both best and median time:

Call	10M best (ms)	10M median (ms)	50M best (ms)	50M median (ms)
argsort int64	7.410 → 7.412	7.555 → 7.563	37.399 → 37.395	38.297 → 38.251
argsort float64, 5% nulls	7.860 → 7.858	8.054 → 8.021	40.840 → 40.753	45.007 → 44.756
argsort utf8	20.402 → 20.339	20.928 → 20.773	120.588 → 121.152	122.698 → 122.940
top_k 10 int64 *	3.898 → 3.811	4.878 → 5.243	6.106 → 6.190	6.915 → 6.960
top_k 5000 float64 *	3.546 → 3.572	4.709 → 4.645	4.491 → 4.494	4.839 → 4.766
partition_nth_indices int64 *	5.592 → 5.554	8.012 → 8.235	11.828 → 11.800	13.081 → 13.070
lexsort (int32, float64)	11.991 → 11.898	12.257 → 12.229	70.462 → 70.462	71.681 → 71.602
rank int32	6.491 → 6.507	6.705 → 6.694	44.901 → 44.798	46.054 → 45.976
group_by keys, 1000 groups	2.660 → 2.687	4.990 → 4.334	4.653 → 4.653	4.798 → 4.772
take int64 through its argsort	1.646 → 1.641	1.704 → 1.702	10.207 → 10.222	10.307 → 10.318
join indices int64, n × n/4	5.103 → 5.098	5.268 → 5.257	44.425 → 44.229	50.299 → 49.478

A median of a call of a few milliseconds moves by several percent from one block to the next (top_k 10 at 10M: 5.453 → 5.012 in the first block, 4.878 → 5.243 alone); a loss shows in both columns.

Chunked columns and the import

A column often arrives as many Arrow arrays: a pyarrow ChunkedArray, a multi-chunk Polars Series, one column of a stream of arrow-rs RecordBatches. am_import_chunks (Sources/ArrowMetal/ChunkedImport.swift) takes all of them in one call and writes each chunk straight into the final page-aligned Metal buffers, with no concatenated copy in between. The work is split by output position, not by chunk, into pieces that the copy threads take one after another, so 6,104 chunks of 8,192 rows and 16 large chunks spread alike, a thread on a slower core takes fewer pieces, and no two threads write the same byte:

• fixed-width values: one memcpy per chunk piece, from the chunk's Arrow offset;
• bitmaps (validity, boolean values): assembled 64 output bits at a time from each chunk's bits at its own bit offset, since chunk boundaries are generally not byte aligned; a chunk without a bitmap, or with a null count of 0, contributes set bits, and the null count is counted from the merged bitmap;
• utf8 / binary and their large variants: offsets rebased by a prefix sum of the chunks' byte lengths and narrowed to int32 (a total of 2 GB or more is refused, as for one large_utf8 array), bytes copied;
• utf8_view / binary_view: views copied with each out-of-line view's buffer index and offset rewritten. Data buffers shared by several chunks (the slices of one array) are taken once; up to 64 distinct buffers are wrapped without a copy, more are copied back to back into merged buffers. A single utf8_view array with more than 64 data buffers takes the same path.

Dictionary, nested, run-end encoded and extension types are not taken: the call returns 3 without reading or moving anything, and the bindings concatenate those columns first, as before.

The view pass. The views are rewritten by am_rewrite_views (Sources/CArrowMetalCopy), C with NEON intrinsics: vld4q_u32 loads four views and splits them into lengths, prefixes, buffer indices and offsets; every lane is rewritten with vector selects (an inline view is kept as it is, an out-of-line one gets its merged buffer index and its offset plus its buffer's base) and checked in the same instructions: length and offset below 2^31, buffer index below the chunk's buffer count, offset plus length within its buffer. vst4q_u32 writes the four back, and the lengths of the valid rows are summed for the column's byte total. Nothing branches per view. A chunk with one data buffer uses its target as a constant; a chunk with several loads each lane's target by its buffer index. A block of 2,048 rows with a view that fails the check is rewritten again view by view, which gives that view an index no buffer has, so it reads as the empty string, as after the single-array import. On one thread (ARROWMETAL_IMPORT_THREADS=1) three 50M-row utf8_view columns in 16 chunks import in 85.5 ms, against 339.3 ms with the view-by-view loop it replaces (6,104 chunks: 134.5 against 381.6 ms; best of three alternating rounds, Benchmarks/results/chunked_import_2026-09-30.csv).

Copy threads. The number of threads is a measured policy (ImportThreads). 1, 2, 3, 4, 6, 8, 12 and 16 threads were timed with Benchmarks/chunked_import_bench.py --threads at 100,000 to 50,000,000 rows (Benchmarks/results/import_threads_2026-09-30.csv), and Benchmarks/import_threads_policy.py turns the sweep into two tables with one rule per import: a thread count is allowed when its CPU time is at most 32 ms or at most twice the one-thread import's, whichever is larger, and the policy takes the fewest threads that come within 5% of the fastest allowed count. Copies have no CPU limit from the size where the limit costs wall time: the smallest measured size from which, at every larger size too, some type and chunk layout is more than 5% slower under the limited count than at its own fastest count. In the sweep that is 381 MiB per column (50M int64 rows in 16 chunks: 4 threads within the limit, 7.0 ms, against 6.5 ms on 16), and the sizes below it, up to 305 MiB, lose nothing to the limit; from there a copy takes the fastest count. Sizes are grouped by the bytes an import writes (both chunk layouts together), and between two measured sizes the boundary is their geometric mean:

copies (values, string bytes, offsets, bitmaps, view data buffers): bytes written	threads
under 1.1 MiB	2
1.1 to 2.7 MiB	1
2.7 to 5.4 MiB	2
5.4 to 108 MiB	8
108 to 341 MiB	12
341 MiB and more (no CPU limit)	16

the utf8_view / binary_view view pass: bytes of views	threads
under 2.4 MiB	4
2.4 to 5.4 MiB	8
5.4 to 10.8 MiB	12
10.8 to 21.6 MiB	8
21.6 to 216 MiB	12
216 MiB and more	8

Below about 2 MiB the copies are split into pieces of at least 1 MiB, so the first row's two threads run one piece; 250,000 int64 rows in 16 chunks take 0.10 ms on 1 to 16 threads. A memory copy of the largest columns stops gaining wall time at four threads until all sixteen cores run it, and those sixteen take 5.6-6.0x the one-thread CPU time in 16 chunks (4.2-4.5x in 6,104 chunks); the view pass, bound by its arithmetic, divides its wall time by the thread count at almost no CPU cost up to six. Wall ms / CPU-ms, three columns, in 16 chunks and (last column) one utf8_view array with 6,104 data buffers, best of 5 runs (Benchmarks/results/import_threads_2026-09-30.csv; the last column is its single_import columns):

threads	int64 10M	int64 50M	utf8 50M	utf8_view 50M	utf8_view 50M, one array
1	3.63 / 3.6	16.5 / 16.5	35.4 / 35.4	85.5 / 85.5	139.3 / 139.3
2	2.52 / 5.0	9.2 / 18.4	19.4 / 38.5	43.3 / 86.4	72.3 / 142.3
4	2.27 / 8.9	7.0 / 27.6	14.3 / 56.2	22.2 / 87.3	41.2 / 158.8
6	2.26 / 13.2	7.0 / 41.1	14.1 / 83.3	15.1 / 88.7	32.1 / 182.2
8	1.46 / 10.7	6.9 / 54.3	14.1 / 110.1	14.3 / 110.7	27.4 / 204.3
12	1.43 / 15.8	6.9 / 80.9	13.9 / 162.9	14.2 / 163.9	27.0 / 298.4
16	1.36 / 19.2	6.5 / 98.3	13.3 / 198.7	13.9 / 204.7	26.3 / 374.2

am_set_import_threads(n) (Python am.set_import_threads, Rust arrowmetal::set_import_threads, Swift ImportThreads.limit) or the environment variable ARROWMETAL_IMPORT_THREADS caps every step at n threads; 1 keeps the import on the calling thread, 0 is the policy.

The page probe. A copy-free import has to know that every page it wraps is mapped. The probe used to be mincore, which reports on every page. mach_vm_region answers per VM region, and a large allocation is one region. Three page-aligned pyarrow int64 columns of 50M rows (400 MB each) imported in 15.71 ms and 15.7 CPU-ms with mincore and import in 0.05 ms with mach_vm_region; at 10M rows 3.10 against 0.02 ms, at 1M rows 0.32 against 0.01 ms (best of two rounds, the single_import columns of the int64 rows in 16 chunks of Benchmarks/results/chunked_import_2026-09-28.csv, where old is the build with mincore). A buffer that is not page aligned is still copied.

In practice pyarrow's buffers of 1M rows and more (pa.array from numpy, compute results, combine_chunks) start on a 16 KiB page and are wrapped; so are arrow-rs concat results of fixed-width and utf8 columns at 1M rows and more.

Measured (M4 Max, Benchmarks/chunked_import_bench.py, three columns per case, the previous build (c4e2f6f, which copied on all sixteen cores) and this one alternating; Benchmarks/results/chunked_import_2026-09-30.csv). At 10M rows: three rounds, best of 5 runs per round, best of the rounds. At 50M rows: each case re-timed alone, eight rounds of 100 runs, best (the 1-minute load was 2.8 when the block started). "combine + import" is combine_chunks() and the import of the combined column (one round); "chunked" is am.array(chunked_array):

3 columns	chunks	combine + import	chunked, previous	chunked	CPU-ms previous / now
int64 10M	16	3.3 ms	1.61 ms	1.54 ms	20 / 11
int64 10M	1,221	4.2 ms	2.80 ms	2.76 ms	20 / 12
utf8 10M	16	6.9 ms	3.37 ms	3.14 ms	40 / 35
utf8 10M	1,221	9.3 ms	4.62 ms	4.42 ms	44 / 37
utf8_view 10M	16	76.3 ms	6.57 ms	3.24 ms	79 / 32
utf8_view 10M	1,221	75.6 ms	11.12 ms	7.89 ms	116 / 65
int64 50M	16	16.5 ms	6.52 ms	6.42 ms	91 / 99
int64 50M	6,104	21.4 ms	14.38 ms	13.88 ms	96 / 106
float64 50M	16	16.5 ms	6.51 ms	6.43 ms	90 / 98
float64 50M	6,104	21.2 ms	14.11 ms	13.74 ms	98 / 108
utf8 50M	16	35.1 ms	13.41 ms	13.21 ms	182 / 198
utf8 50M	6,104	47.6 ms	22.15 ms	21.27 ms	209 / 227
utf8_view 50M	16	372.6 ms	31.19 ms	14.26 ms	392 / 111
utf8_view 50M	6,104	379.3 ms	54.26 ms	38.96 ms	564 / 299

Every case in the table is as fast as on the previous build or faster on the best run, and on the median run every case but utf8 at 50M in 16 chunks, whose median run is 14.77-14.82 ms in each of the eight rounds against 14.21-14.45 ms (Benchmarks/results/chunked_import_2026-09-30.csv, as are the figures below). CPU time: the 50M copies, where the policy has no CPU limit, take 1.09-1.11x the previous build's (and 4.8-6.0x the single-threaded combine_chunks and import, for 0.38-0.65x its wall time); the 10M copies take 0.54-0.88x; the utf8_view pass takes 0.28x at 50M in 16 chunks and 0.41-0.56x elsewhere in the table. At 20M and 30M rows (three rounds) every copy case is within 1.3% of the previous build's best time or faster, at 0.81-1.22x its CPU time (utf8 at 30M in 16 chunks: 8.67 against 8.56 ms, 128 against 104 CPU-ms); at 1M rows within 0.01 ms or faster (utf8_view 0.52 against 0.95 ms in 16 chunks). utf8_view imports 1.08x the time of utf8 in 16 chunks (14.26 against 13.21 ms) and 1.83x in 6,104 chunks (38.96 against 21.27 ms): there its 6,104 data buffers are copied into one merged buffer, so the column moves 16 bytes of view per row plus the string bytes, against 4 bytes of offset per row plus the string bytes for utf8.

One utf8_view array with 6,104 data buffers (combine_chunks of the 6,104-chunk column above) imports in 26.5 ms and 274 CPU-ms (previous build: 44.8 ms and 548 CPU-ms; the single_import columns of Benchmarks/results/chunked_import_2026-09-30.csv; the build before the merged copy path: 150.5 ms and 143 CPU-ms, Benchmarks/results/chunked_import_2026-09-28.csv); at 10M rows in 5.6 ms and 59 CPU-ms against 9.2 ms and 112 CPU-ms (chunked_import_2026-09-30.csv).

Through the engines, with the previous build alternating: Benchmarks/polars_engine_bench.py --sizes 2000000,50000000 --cases a,c,i,t1,w1,x2, MetalEngine with the import cache cleared, takes 0.98-1.01x of the previous build's best time at 50M rows (three rounds) and 0.73-1.06x at 2M rows (the 1.06x, the 200-group sum under the default engine, takes 0.91x of its median); the two 2M cases that differed more with every shape taken, the 100,000-group sum and the join, re-timed alone over eight rounds of 100 runs, take 0.92x and 0.98x of its best time and 1.02x and 1.00x of its median (Benchmarks/results/chunked_import_polars_2026-09-30.csv). The DataFusion rule's sorts (datafusion/examples/bench.rs --families sort, rule on, two rounds, datafusion/results/datafusion_sort_import_2026-09-30.csv) take 0.93-0.99x of the previous build's time at 50M rows in both table layouts, at 0.91-1.04x its CPU time (int64 key in 8,192-row batches: 68.84 against 70.63 ms), and 0.95-1.05x at 10M rows with 0.54-0.88x its CPU time.

• Not a query planner or SQL engine. It is the compute layer that DuckDB, DataFusion, Polars plugins or an app can call.
• Not a tensor library; MLX is that.