Fastest Aho-Corasick on GPUs - #334
Open
ashvardanian wants to merge 50 commits into
Open
Conversation
Introduces the `find_many` headers: a goto-completed Aho-Corasick automaton built once on the host and matched against a whole batch of haystacks, by CPU engines and CUDA kernels sharing one published `aho_corasick_view`. Transitions live in two tiers - a dense 256-wide table for the most-visited states, a double array with failure links for the rest - so the split follows how often a state is actually visited rather than the alphabet. Matches are reported through a CSR output pool that inherits along failure chains, and `find_many_uncased_k` bakes full Unicode case folding into the transitions themselves via preimage reconvergence, so the haystack is never folded while matching. The GPU backend gives one thread one chunk of a concatenated tape, warming up `max_match_bytes - 1` bytes ahead of its chunk so chunking stays exact rather than approximate, and stages a frequency-ordered prefix of the hot tier into shared memory. A single kernel serves both the counting and scattering pass, selected by a compile-time pass enum. State ids are parameterized over `u16` and `u32` to keep occupancy high, while every container index and byte offset is `size_t`, so one haystack may exceed 4 GB. `scratch_amount_t` moves up from the similarities engine into the shared types header, since more than one family now lays out scratch that way.
Expose the Aho-Corasick multi-pattern engine through the stable C ABI: `szs_find_many_init` compiles a dictionary from a needle collection with cased or case-folded matching, and `szs_find_many_count` / `szs_find_many` run over `sz_sequence_t` and both tape shapes on whichever device scope the caller names. The backends type is a `std::variant` of bare engines selected by the capability mask at construction, dispatched by the same capability-driven visitor every sibling domain uses. Normalize the engine signatures across all three backends: `try_build` takes needles plus case sensitivity everywhere, and `try_count` / `try_find` share one shape - per-haystack counts with a grand total, or a match span with a written count - on serial, parallel, and CUDA alike. One offset-based `find_many_match_t`, layout-identical to the C struct, flows from the dictionary callback to the caller's array with no conversions or staging copies on any CPU path; the CUDA engine owns its host dictionary, uploads to managed memory at build, and drains host-bound outputs with one stream-ordered driver copy. Harden the width discipline along the way: the cold-tier probe indexes wide before narrowing, edge ordinals gain a ceiling check against the shared `state_id_t` bound, and the parallel scatter guards its overlap pointer on empty dictionaries. Cover it with an adversarial needle-by-haystack test sweep against a brute-force oracle, backend-agreement and buffer-contract checks, a forced all-cores-on-one-haystack slicing test, device-memory contract tests for scattered unified allocations and host-input rejection, and a Zipf-sliced benchmark reporting construction in needle-bytes and search in haystack-bytes per second across vocabulary cutoffs.
The multi-pattern family is growing operations beyond locating matches - `replace` rewrites and BM25-style scoring share the same compiled Aho-Corasick dictionary - so it takes a plural product noun beside `similarities` and `fingerprints` instead of one operation's verb. The bare locate call `szs_find_many` becomes `szs_substrings_find`, and its tape twins follow the operation name as `szs_substrings_find_u32tape` and `szs_substrings_find_u64tape`. Every other symbol renames mechanically, and `aho_corasick_*` keeps its algorithm name.
Nothing in the tree ever filled the struct, so every engine tuned to a fictional 256 KB L2 and 8 MB L3. The spawned `forkunion_executor_t` now answers with a `specs()` method reading its own detected topology: the deepest domain-confined cache fills `l3_bytes` - 105 MB instead of 8 MB on this box - and the logical-core and memory-domain counts fill the core fields. `l1_bytes` and `l2_bytes` keep their conservative defaults, since the ForkUnion C API exposes no per-core cache-level query yet. The C runtime's CPU scope stores the spawned executor's real specs instead of a defaulted local, so every `get_specs` consumer across the five domain shims inherits the fix. The serial default scope keeps defaulted specs - it owns no topology to ask. All three benches now thread `pool.specs()` into every CPU engine cell beside the executor, mirroring how the GPU cells already pass their fetched `gpu_specs_t`. The UTF-8 baselines spell the `dummy_executor_t` explicitly to reach the specs slot of the greedy executor overload. Tests intentionally stay on hand-crafted specs - they force both sides of every cache threshold deterministically on any machine.
Five cross-TU helpers in the CPython extension were defined with `SZ_API_RUNTIME` while their declarations in `python/stringzilla/stringzilla.h` are plain `extern`. On MSVC that macro is `__declspec(dllexport)`, so the pair is a linkage mismatch and both Windows jobs failed with `error C2375: redefinition; different linkage`; GCC and Clang accept it silently. Nothing inside the module needs exporting beyond its init function, and the sibling helpers in the same files - `sz_py_is_mutable`, `sz_py_export_optional_index`, `Strs_get_start_` - were already plain.
`sz_assert_` is gated on `SZ_DEBUG`, and `CMakeLists.txt` defines `SZ_DEBUG=0` for every configuration but `Debug`, where `types.h` expands it to `((void)(condition))`. Every CI test build is `Release` or `RelWithDebInfo`, so 77 assertions in the GPU test families evaluated their argument and discarded the verdict - the differential comparison at the heart of the similarity suite could not fail. The rest of `test/` had already moved to `verify()`, which no configuration strips; these two files and one helper in the shared header had not. 24 of them performed the work inside the assertion, such as filling a tape through `try_append`. Those become `let_verify`, hoisting the call into its own statement so the result is named before it is checked, rather than depending on `((void)(condition))` still evaluating its argument. Also drops the `unique` parameter of `randomize_strings`: no caller has ever passed it, and the path was unsound - it deduplicated after `resize`, handing back fewer strings than the requested batch size with no diagnostic.
A helper that chooses or encodes now returns its result and the caller appends, so every arm of a corpus `switch` reads the same way instead of mixing `append_codepoint_(out, cp)` against arms that append inline. `append_codepoint_` becomes `encoded_rune_`, returning a short-string-optimized `std::string` that never reaches the allocator at four bytes, across 53 call sites. The builders that stream many items into a reused buffer keep their sink parameter - the rule is about choosing and encoding, not accumulating. `utf8_runes.cpp` and `utf8_tokens.cpp` included `stringzilla.hpp` directly and re-derived helpers the shared header already had, so `append_codepoint_` was triplicated and `random_valid_utf8_` duplicated verbatim. Both now include `utf8.hpp`, which gains `random_valid_utf8_` and `random_valid_utf8_bytes_` as their single home; the hex-dump helper collapses from four copies under two names to one. Spells out `rng` as `generator` in the 16 signatures that carried it, matching the call sites, and gives the two sentence straddle builders the family infix every sibling already had.
The adversarial sweep now crosses three declared axes - needle shapers, byte-agnostic placement skeletons, and needle transforms - instead of two fused generator enums, and every planting records its ground truth: the exact match record and whether the engine must report or must miss it. A new referee asserts those declared effects on every cell, a third witness beside the brute-force oracle that cannot share a misconception with either walk, since the generator knows where it planted what. Vocabularies now sample the shipped fold-preimage tables instead of hand-picked literals: narrow-image needles cover Greek, Cyrillic, Cherokee, and astral fold pairs by construction, wide-image needles exercise 2:1 and 3:1 byte expansions, and a mixed-width shaper drags every needle through all four UTF-8 decode widths. One table-driven inverter re-spells needles through fold space for the fixtures and the agreement pool alike. The old near-miss arm was an ASCII case flip mislabeled as a negative control - under an uncased dictionary those haystacks still match. The transform axis states each effect per sensitivity mode, and a true negative control perturbs a non-case bit. New placements interrupt variants mid-codepoint with an unbridgeable byte, and plant matches across slice seams of one large haystack that sliced-L2 cells walk with a real thread pool; cased cells also sprinkle malformed continuation bytes into the noise. The independent oracle returns matched haystack slices as spans rather than offset pairs, and both consumers recover the record fields in place.
Add STRINGWARS_DATASET_LIMIT to cap how much of the corpus each benchmark reads, chosen by what the benchmark measures. Compute-bound families take a 64 MiB slice, enough to exercise every control-flow path, while memory-bound families read the whole file so their steady-state bandwidth is measured on as much data as is available. The bound is applied at read time, so a capped run never touches the file tail. This replaces STRINGWARS_MAX_TOKENS, whose token count conflated a byte budget with a byte-independent token budget.
Aho-Corasick walks a data-dependent transition load, so the CUDA engine is latency-bound rather than bandwidth-bound - it sustains under one percent of an H100's memory throughput. Resident warps are the only thing hiding that latency, which inverts the staging trade: the hot tier is now copied into shared memory only when all of it fits without costing a resident block, and read through the cache hierarchy otherwise. A partial prefix pays the copy per block and a bounds test per byte while most transitions miss it anyway, so staging is all-or-nothing. The budget behind that rule comes from the occupancy query rather than from arithmetic or a tuned constant. Dividing a multiprocessor's shared memory by a target block count silently yields one block fewer, because the driver charges every block a reserve of its own and rounds to a granularity no attribute reports; `cuOccupancyAvailableDynamicSMemPerBlock` repeats that same division, so `shared_memory_budget_for_resident_blocks` binary-searches `cuOccupancyMaxActiveBlocksPerMultiprocessor` instead, and a zero target derives its block count from the occupancy the kernel reaches with no shared memory at all. Measured on an H100 over 64 MiB of news text, counting rises from 14.8 to 29.2 GB/s and finding from 8.1 to 28.2 GB/s at 3.4K needles, and from 5.2 and 2.5 to 14.8 and 10.2 GB/s at 346K needles, with match counts unchanged throughout. The dictionary now publishes a walking automaton rather than the raw spelling trie. A case-folded trie is a DAG - several spellings reconverge on one state at different byte widths - so its failure links are not single-valued. Splitting each state by its `(spelling, failure)` pair is the coarsest split that restores them. The doubling family, "ssssss" folded, is provably exponential and is declined with `overflow_risk_k` once the walking count would overflow the id type, rather than built wrong. Both walkers feed four transitions per tape load. The state chain stays strictly serial - only the fetch widens - so `count` and the parallel prefix counter drop three of every four loads. On the device the cursor is peeled to its own 4-byte boundary and the load issued as one `ld.u32`: PTX traps on a misaligned `ld.u32` rather than slowing down, so the compiler otherwise lowers a 4-byte `memcpy` from an unproven pointer into four byte loads plus three `prmt` merges. The CUDA walk tests one acceptance bit before touching `outputs_counts`, so the common no-match byte never reads the global counts array. The warm-up prefix is peeled out of the emit loop instead of re-testing its bound on every byte, and both passes share one `chunk_match_slots` buffer, which the host's in-place exclusive scan had already made the same allocation. Hot-tier ordering is shallow-first with each depth band sorted by out-degree descending. Per-state growth during insertion reserves in power-of-two steps, since `try_reserve` allocates exactly what it is asked and a bare `try_resize(count + 1)` per state re-moved the whole array every time.
…chmarks Every StringZillas test and benchmark inherits `SZ_DYNAMIC_DISPATCH=1` from the engine libraries it links, which externs the `sz_*` entry points and leaves them to be resolved at link time. Nothing resolved them, so the moment a benchmark called one directly - `sz::string_view::split` and the `sz::hash` / `sz::equal_to` pair behind the substrings vocabulary - both bench targets failed to link. Neither shipped library could take that job: `stringzilla_shared` and `stringzilla_bare` both set `SZ_OVERRIDE_LIBC=1`, interpose LibC's own symbols, and fault the process before `main`. The shim units default `SZ_OVERRIDE_LIBC` to `SZ_AVOID_LIBC`, so `stringzilla_static` leaves both at zero and is otherwise the same dispatch core, built only when tests or benchmarks are. It carries no install rule, being an in-tree consumer's archive rather than a shipped one. `define_stringzilla_shared` becomes `define_stringzilla_library`, taking the link type as a parameter, so the new target reuses the aliases, position-independence, baseline architecture flags, and SIMD definitions instead of copying them. The substrings vocabulary now counts word frequencies over the whole dataset rather than over `env.tokens`, so `lines` and `file` searches draw needles from the same vocabulary a `words` search does, and `STRINGWARS_UNIQUE` no longer silently flattens every count to one. One hashed pass replaces sorting every occurrence, since the corpus holds hundreds of millions of words but only a few million distinct ones. Words under 3 bytes match at nearly every position and would measure match materialization rather than the walk; words over 32 bytes are unsegmented CJK or Thai runs and URLs, and the longest needle sets `max_match_bytes`, which every GPU chunk re-walks as its warm-up. The README tables are measured rather than pending, in MB/s per cell with the fastest backend of each column in bold, on 64 MiB of `xlsum.csv` split into lines.
Engines that size a working set against the cache the GPU actually reads through had nothing to read it from: `gpu_specs_t` carried VRAM, constant, and shared memory but not L2, so anything cache-resident had to borrow `cpu_specs_t` and inherit a host cache size instead. `CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE` fills it alongside the attributes `gpu_specs_fetch` already reads in its one cold pass. The default stays an A100's 40 MB, matching how every other field in the struct documents its fallback.
…d capacity The CUDA engine chose its hot/cold split from a default `cpu_specs_t`, so the tier a GPU walks was sized by an 8 MB host last-level cache regardless of the device: 8192 rows on an H100 whose own L2 holds five times that. It now sizes against `gpu_specs_t::l2_bytes`, the cache the walk actually reads through, which puts the whole automaton in the hot tier for small dictionaries and raises counting from 29.2 to 30.3 GB/s and finding from 28.2 to 31.3 GB/s at 3.4K needles, and counting from 14.8 to 17.0 GB/s at 346K, on 64 MiB of news text. A refused output buffer no longer zeroes the count it just computed. Every `try_find` counts internally before it can check capacity, so the number the caller needs was already in hand and thrown away, leaving a size query to recover it with a second walk - and on the GPU a second counting kernel, a second device-wide scan, and a second blocking synchronize. The status already separates a refusal from a success, so reporting the need costs no ambiguity and makes an empty buffer a single-call size query. The uncased fixtures that pin the walking automaton's reconvergence splits come with it: the codepoint-boundary end test in the independent oracle, the `(spelling, failure)` hazard table, and the doubling family declining with `overflow_risk_k` at the `u16` ceiling while `u32` still compiles it.
…er state id A `u16` automaton halves the transition row - 512 bytes against `u32`'s 1024 - so twice as much of it stays cache-resident, which is what the hot tier is sized against. Reaching that width meant guessing it before construction, and the guess cannot be made: for cased vocabularies the trie node count is exactly the sorted needles' summed prefix deltas, but folding turns the uncased automaton into a DAG whose `(spelling, failure)` split multiplies states unpredictably. So the width stops being an input. `try_build` gains an overload taking an already-built wider dictionary and adopting its published arrays element by element, which costs a copy rather than a second derivation of the walking automaton. Every ceiling the narrower id imposes - published slots, needle index, match length, outputs per state - is tested before anything is copied, and a vocabulary that does not fit declines with `overflow_risk_k` leaving the wider dictionary usable. It is named by the dictionary rather than the engine holding it: the narrowing needs nothing else, and the concrete parameter type is what keeps the overload from competing with the needles one, which an unconstrained template argument could not. Measured on an H100 over 64 MiB of news text, counting a 3,462-needle vocabulary of 19,291 states rises from 30.9 to 34.0 GB/s once narrowed, with identical occurrence counts; a 34,620-needle vocabulary of 144,929 states declines as it should. Both state-id widths of both CPU capabilities are now instantiated in the shim's translation unit. The serial engines were absent, so every consumer recompiled them despite the C shim building one for a single-threaded scope.
Matching gains a per-call overlap policy: every match as before, or a leftmost cover resolved by longest span or by lowest needle index. One dictionary serves all three, since the policy never shapes the compiled automaton. It is resolved in exactly one place - `count` and `visit` pick the walk - so neither engine branches on it and a new backend implements that pair and nothing else. The overlapping tally keeps its dedicated four-byte-load walk, which is a genuinely different kernel rather than a duplicated dispatch. Rewriting substitutes each match with its needle's replacement, tape in and tape out, because a rewrite's product is itself a tape. It sizes with a zero-capacity call and refuses a short buffer whole while naming the size it wanted, so the capacity contract matches the one `find` already offers. An overlapping rewrite is refused outright: two matches sharing a byte have no single answer, so it is not a function. Scoring walks each haystack once for classic BM25 against one weight per needle. The dictionary is the query here rather than an index to probe, so a second weighting is a second scan, not a second row of a matrix that would be mostly zeros over a large vocabulary. Term frequencies are raw counts and take no policy - a leftmost cover suppresses genuine occurrences of any needle nested in another, which is no longer a term frequency - and the reduction runs in ascending needle order, so scores are bit-stable across runs of one backend. The leftmost walk settles a start once the scan is `max_match_bytes` past it, holding the undecided ones in a window rounded up to a power of two so the slot lookup masks instead of dividing by a value known only at runtime. It drains to the last start a match actually claimed rather than to the end of the haystack, which is what keeps a haystack no needle hits from paying a slot visit per byte. The C entry points for scoring and rewriting land separately.
`sz_capabilities_to_string_implementation_` wrote into a function-local `static char[256]` and returned a pointer to it, so two threads formatting capabilities at once overwrote each other's text, and so did two calls whose results were meant to be compared. It now takes the caller's buffer and its capacity, truncates rather than overflowing, and returns the length written. The four Python call sites pass a stack array. The exported `sz_capabilities_to_string` keeps a `static`, because its signature returns a string it never receives - that is now the only shared buffer left.
`SZ_HELPER_AUTO` now carries `constexpr` wherever the language has the word, so every scalar helper folds at compile time when its arguments allow it - and, more usefully, becomes callable from a CUDA kernel: `nvcc` reaches a host `constexpr` function from device code under `--expt-relaxed-constexpr`. The C layer therefore never has to name an execution space of its own, and needs no `__host__ __device__` token, no per-function portability macro, and no include-order rule. Three constructs stood in the way, all of which C++23 legalizes and no `nvcc` accepts - CUDA 13.2 still tops out at `-std=c++20`: - The uncased danger-zone search chained three near-identical blocks with 13 `goto`s into 3 labels, one per folded rune the needle's anchor could sit on. They collapse into a loop over the anchor index, calling a single `sz_utf8_uncased_verify_at_folded_rune_`. The runes of one source codepoint on either side of the anchor never reach a folded iterator, so they are compared against the image directly, which is what the per-block pre-checks were doing by hand. Behaviour is identical across 1184 fold-anomaly cases, 533 of them matches, at a measured 2-5% cost on that path alone - it fires only on chunks holding one of the 139 boundary-breaking codepoints. - Two AVX2 left-pack tables were function-local statics; they move to file scope unchanged. The folded rune iterators move to `utf8_uncased_fold/serial.h`, beside the fold they call, which is also the header a GPU walk can include on its own. `-Wno-invalid-constexpr` joins the C++ warning set: a helper reaching an intrinsic can never be constant-evaluated, which the diagnostic correctly reports and which is not a defect here.
The C shims carried three guard schemes at once - `STRINGZILLAS_SZS_<FILE>_CUH_` for five files, `STRINGZILLAS_RUNTIME_CUH_` for a sixth, and `STRINGZILLAS_SCAFFOLDING_CUH_` for a seventh, which named a concept rather than its file. They now mirror the path under `c/` behind the prefix every function in them already uses: `SZS_SUBSTRINGS_CUH_`, `SZ_DISPATCH_H_`. Headers under `include/` already mirrored their paths and are untouched, as are the three umbrella guards a downstream consumer might test for. 44 `@file` annotations pointed at paths that have not existed for a long time - most of them at a retired `scripts/` directory with `bench_`/`test_` prefixes, the rest at bare filenames or a renamed table header.
The multi-pattern engine and `sz_utf8_uncased_search` disagreed about what a case-insensitive match is, and the multi-pattern side was wrong. It inverted the fold onto the needle - enumerating every source codepoint whose fold equals a 1-3 rune window and giving each a raw-byte trie path - so every match had to consume whole source codepoints, and `rune_arrival_t::atomic_k` existed to document that refusal. A needle "s" could not match inside "ss". That refusal was not a bug to patch but the only semantics the architecture could express. Any deterministic automaton whose alphabet is raw haystack bytes and which must report match spans is exponential in needle length: the fooling set is a run of sharp-S against a run of long-S pairs, fold-identical at 2 bytes versus 4, so the reported span encodes the byte width of every spelling behind it and all prefixes are pairwise inequivalent. Uncased "ssssss" measured 454824 states, and wider dictionaries declined with `overflow_risk_k`. The fold now happens in the byte stream instead. `substrings_folded_cursor_t` walks one codepoint at a time and feeds folded bytes to a plain trie, so a needle contributes one path and the automaton is the cased one. Matching is the fold-subset rule the single-pattern kernel already follows: `fold(needle)` is a contiguous run of `fold(haystack)`, with both ends snapped outward to codepoint boundaries. "ssssss" is now 7 states. Folded and source offsets agree everywhere except within a window touching one of the 139 boundary-breaking codepoints, so `substrings_folded_span` resolves a match by subtraction in the common case and walks backwards only past the last break. `substrings_resolve_match` also settles whether an earlier rune of the same codepoint already reported that span, which is what keeps "s" in a sharp-S to one match while leaving "ss" over two sharp-S at three distinct spans. `sz_utf8_lead_may_fold_` answers "can anything under this lead byte fold" from four immediates, with no table and no memory access on either a CPU or a GPU; without it, deleting the old classifier cost 27% on Chinese and 23% on Arabic. `match_bytes` splits into folded and source pairs throughout: the automaton walks folded bytes, every haystack extent - halo widths, slice sizes, pending starts - needs the source bound, and the two must never be compared. Validated against an independent fold-subset oracle over 3000 cased and 30000 uncased random dictionaries with zero mismatches, plus per-needle agreement with `sz_utf8_uncased_search` itself.
The device walk stepped raw haystack bytes, so an uncased dictionary would have answered differently from every CPU backend. It now runs the same `substrings_folded_cursor_t` the host walk uses, resolving each match through the shared `substrings_folded_span` - one subtraction while folded and source offsets agree, a bounded backward walk past the last boundary break. Leftmost policies reach the GPU with no host resolution and no atomics on the hot path. A survivor whose start lies at or after the maximum end seen so far is accepted unconditionally, independent of everything before it, which makes the cover resolvable inside one chunk: each thread keeps a small ring of pending starts and settles them locally. Measured over the corpora, the prefix-max segments have a median length of 1 and a maximum of 29, all within the ring. That ring is `substrings_device_pending_width_k` wide, which nothing previously checked against the dictionary. A needle whose worst source span exceeds it would have aliased slots silently, so `try_count`/`try_find` now refuse with `overflow_risk_k` up front. The leftmost walk also gates on the acceptance bitmap its two sibling walks already used, instead of discarding it. `bench/substrings.cuh` had never compiled: its `try_count`/`try_find` calls omitted the `overlap_policy` argument the engine has required since before this work.
`multiplying_rolling_hasher` carries no modulo and says so - it relies on the integer type's natural overflow. That is defined for unsigned types and undefined for signed ones, so instantiating it over `i32_t` made every `push`, every `roll`, and the `highest_power_` loop undefined at exactly the width they were meant to wrap at. It is not a sanitizer nuisance: a compiler may assume signed overflow cannot happen and optimize on that basis, and the reported values sat right at the signed maximum. `buz_rolling_hasher` needs the same guarantee for a different reason. Its rotation shifts both ways, and a signed left shift pushes a sign bit out while a signed right shift replicates it rather than feeding in zeros, so the rotation would neither be a rotation nor be defined. Only unsigned instantiations exist today, and now only those compile. `rabin_karp_rolling_hasher` already named the three widths it accepts and needed nothing. Signedness buys nothing here to weigh against that. Two's complement makes addition, subtraction, and the low half of a multiply the same operation either way, and `push`, `roll`, and a whole-buffer sweep compile to identical instructions on x86-64 and AArch64 alike. Signed arithmetic does help a loop induction variable, where undefined overflow lets an index widen into 64-bit addressing, but that is a different variable from the accumulator. At the point of use it costs instead: reducing a hash modulo a slot count is one mask for an unsigned value and a sign correction for a signed one. So the state types are constrained rather than the arithmetic widened, which keeps the hot path exactly as it compiles today. The narrow widths deserve the same scrutiny for a different reason - anything below `int` promotes on every operation and smuggles the same signed overflow back in - but no such instantiation exists.
Three ways the build silently did the wrong thing with CUDA, all of them independent of which compilers a machine happens to carry. `check_language(CUDA)` caches the compiler it finds together with the host compiler it happened to pair with it, and that pairing then outranks a caller's `CMAKE_CUDA_HOST_COMPILER`, so setting one was ignored without a word. The probe now runs only to pick the default for the option; asking for CUDA outright leaves the detection to `enable_language`. `-march=native` for CUDA sources was decided by `check_cxx_compiler_flag`, which puts the question to the C++ compiler - the wrong toolchain. NVCC delegates host compilation but parses the host's headers itself on the device pass, so the pair decides, and the C++ compiler accepts flags NVCC then chokes on. A probe over `probes/cuda_native_arch.cu` asks NVCC instead, carrying the header that actually settles it. `-Wno-invalid-constexpr` never reached the CUDA host, where a Clang host makes that diagnostic an error and every source including the `constexpr` C layer fails. The presets carry intent and no versions: which host compiler a toolkit accepts is a fact about the machine, so it belongs in the per-developer `CMakeUserPresets.json`, now ignored.
The GPU backend gained the three operations it was missing. A leftmost cover now resolves on-device: every chunk emits its matches, a resolve pass marks which survive, and a compaction keeps them - so the walk is never repeated to decide ownership. On top of that, `try_replace` rewrites whole haystacks on the device, splicing each match's replacement against offsets a scan produced, and `try_score_bm25` reduces each haystack to one float, merging per-block frequency rows in a shape no grid size perturbs. A leftmost `try_find`, a rewrite, and a scoring pass no longer fall back to the host. On the CPU, a haystack too large for one core's cache is now rewritten by every core at once. Each core owns the matches starting in its slice and writes into a disjoint output range, so neither a mutex nor an atomic sits on the write path, and the per-core shares settled while sizing survive into the writing pass - the cover is resolved once per haystack rather than twice. Construction splits in two. `szs_substrings_init` takes only the capability mask, like every sibling engine, and `szs_substrings_index` compiles a needle set against a named device. The hot/cold tier is sized to fill the cache that walks it, so the tuning input and the tuning decision now sit in one call; indexing again replaces the needle set, which is what makes re-tiering for another device something a caller asks for rather than something that happens silently. The state-id width stops being an engine template parameter: each engine holds whichever of the two automata its needle set fits, instead of leaking u16/u32 out to the C shim as two engine types and a four-arm backend variant. The CUDA engine also stored its automaton twice. Its allocator is unified memory, reachable from every device, so the seven `device_*` mirrors and the upload that filled them were copies from unified memory into unified memory; there is no `cudaMemcpy` anywhere in the file. The kernels now walk the dictionary's own arrays, and the kernel table carries one shape per state-id width in a single per-device cache. `szs_substrings_stats` is dropped along with the struct it filled.
The two chunk-walk launchers built byte-identical ten-argument blocks, differing only in which kernel handle they passed and whether the output span was empty - which is exactly what `substrings_pass_t` already names. They collapse into one `launch_walk_at_`, so the argument order the kernels depend on is written once rather than three times, and the docstring claiming as much stops being wrong. The three per-width shape accessors were the same ternary three times over; `kernels_t::by_width_t` now answers for itself. Three of the parallel engine's scratch buffers handed out rows through a named accessor and three did the offset arithmetic inline, at six call sites. Every row is now reached the same way, which makes the BM25 merge legible for the first time: it indexes the touched row by slot and the frequency row by needle index, where the old flat arithmetic hid both behind one shared offset. Release codegen is unchanged - the accessors fold to the same pointer arithmetic. `try_reserve_rewrite_shares_` sized two unrelated buffers under two separate guards and spent eight docstring lines justifying the pairing; the coverage row gets its own reserve and the paragraph goes. The benchmark spelled out a four-operation by three-backend grid cell by cell, about twelve repetitions of allocate-wrap-name-measure. The backends differ only in which engine runs and in the `(executor, specs)` pair every entry point already takes, so one shape per operation covers all three. Skip conditions stay where they were: the occurrence cap, GPU availability and the rewritable policy are control flow, not repetition.
`try_index` was the only substrings entry point whose executor template parameter carried no `executor_like` constraint, so passing specs in the executor slot compiled. Three benchmark call sites did exactly that, and the automaton was tiered against a default one-core machine rather than the pool's probed cache - the benchmark reported hot and cold tier sizes it had not measured. Constraining the parameter like its four siblings turns that into a compile error. The CUDA twins of those calls passed a `gpu_specs_t` into a `cuda_executor_t const &`, which cannot bind at all. That they were committed is direct evidence the CUDA substrings paths were never built after the signature changed; the GPU rewriting and scoring tests had never compiled. They do now, and pass. The benchmark's timed callables capture their executor by reference and outlive the expression that builds them, so the executors they were handed as temporaries died before the first measured iteration. Both are now owned for the whole sweep cell and passed as lvalues. Two test sites dropped a probed `gpu_specs_t` that was live in the same function while every neighbouring call threaded it.
Every batch operation in the C layer takes three input shapes - `sz_sequence_u32tape_t`, `sz_sequence_u64tape_t`, and the callback-addressed `sz_sequence_t` - but `AnyBytesTape` and `AnyCharsTape` carried only the tapes, so the sequence entry points were reachable from the owned-return `compute` and from nowhere else. A caller holding `Vec<String>` had to materialize a tape before `compute_into` would look at it. Both enums gain the missing arm, built by `from_slices`, and every `compute_into` dispatches it to the sequence extern it already declared. `SzSequence` becomes a public opaque type carrying the lifetime of what it borrows, which is what keeps the raw handle honest, and sheds the two fields it had beyond the four `sz_sequence_t` actually defines. Mixing a tape with borrowed slices across the two sides of a cross-product stays refused: satisfying it would mean copying one of them behind the caller's back.
`szs::Substrings` compiles a needle set into one Aho-Corasick automaton and reuses it across every later call: `count_into`, `find_into`, `score_bm25_into`, `replace_bound` and `replace_into`, each writing into caller-owned buffers so a pipeline allocates once. Haystacks arrive as `AnyBytesTape`, so all three input shapes the C layer accepts reach the engine; only the rewrite narrows to the tapes, since its product is itself a tape and a callback-addressed sequence has nowhere to put one. `Bm25Params` has no `Default`. A corpus mean is a property of the corpus rather than a tunable, and pairing the literature's `b = 0.75` with a zero mean asks the engine to divide by a mean that is not there. The two configurations that exist get names instead - `normalized(mean)` and `unnormalized()` - so the impossible one is unrepresentable. `aho-corasick` arrives as a dev-dependency, not an optional feature, so it never reaches a consumer's build graph. Its three `MatchKind`s are our three overlap policies under different names, which makes it an independent witness for every policy and for the rewrite; the tests differential all four against it over seeded corpora.
`szs.Substrings` exposes the engine through a `tp_methods` table rather than the single `tp_call` its siblings use, since `count`, `find`, `score_bm25`, `replace` and `replace_bound` cannot share one keyword set. `find` returns four `uint64` arrays instead of one structured array, so a caller wanting a single column takes it without learning a compound dtype, and it negotiates its own capacity through the zero-capacity size query the C layer documents. `sz_py_export_strings_as_sequence` becomes total over all five `Strs` layouts. Needles have no tape overload, so this is the first caller that reaches it without probing the tapes first, and its `STRS_FRAGMENTED` assertion was a precondition only the probe order upheld. `average_document_length` is required rather than defaulted: it is a property of the corpus, and a zero mean beside the customary `b = 0.75` asks the engine to normalize by a mean that is not there. The rewrite answers an empty batch before it sorts the input into a layout. An empty `Strs` is fragmented when built from a list but becomes a tape whose offsets are still NULL once the unified allocator swap runs, so reading the trailing offset dereferenced NULL on a CUDA engine and refused the same input as "not a tape" everywhere else. `stringzillas` links the core sources for the first time here: the rewrite path calls the dispatched `sz_copy`, the way every CMake target links `stringzilla_static` rather than leaving it undefined.
…uites BM25's `average_document_length` was a sentinel: a zero mean collapsed the length term to one, so it silently discarded both `length_normalization` and every `document_lengths` entry the caller had computed - and said what `length_normalization = 0` already says. The C entry point now refuses the pair, once, for all three languages, and the header states what the kernel's own floor is for. The C++ binaries register the four substrings tests that were declared but never run, and the `cuda_clang` preset pins all three compilers rather than the CUDA host alone. Clang mangles a C++20 requires-clause into the symbol name and GCC does not, so compiling the C++ sources with one and the CUDA sources with the other gives the constrained engine entry points two different symbols and no link. Two guards join the pre-commit hook: one for `clang-format`'s `BreakStringLiterals` mangle, whose signature is two literal fragments left on one line, and one for indexing straight through a call, where the temporary the index reads from dies at the semicolon.
All three `substrings_pass_t` enumerators read as "count something", and the one that actually counts occurrences per needle was the one not called counting. Rotate them so the name states the pass: `sizing_k` sizes the output, `writing_k` writes matches at known offsets, `counting_k` counts. `shared_rows_` becomes `staged_rows_` to pair with the `staged_accepts_words_` it is budgeted and launched beside - one concept that carried two prefixes.
…8 battery The C++ `stringzillas` suites never ran in CI. `test_ubuntu_cpus` already built `stringzillas_test_cpp20` on every push and then discarded it, so 14 substrings, 5 similarities and 3 fingerprints tests were only ever exercised by hand. It now runs there, which costs runtime alone since the binary was already compiled. Five families hand-rolled the same malformed-input sweep that `for_each_adversarial_utf8_input_` already drives for the segmentation tests. They now share it. `uncased` gains the most: it walked all 256 singles and all 65,536 pairs regardless of `SZ_TESTS_MULTIPLIER`, never replayed through `for_each_cacheline_offset_`, and skipped the four named malformed shapes - so sharing the battery makes it both cheaper and stricter, and it finally honours the dial. `test_similarities_memory_usage` documented its table as cheapest-first so a reduced multiplier keeps the cheap prefix, but sized itself from the table length, so the default pass ran every row including the 8192-byte ones, and a raised multiplier resized past the table and appended default-constructed experiments rather than harder ones. Both ends are fixed: 10.4s to 0.3s at the default, still widening to the full table at 2x. Device tests are gated inside their bodies rather than around their definitions, so `stringzillas.cpp` and `stringzillas.cu` register identical lists and cannot drift apart again. The repeated unified-memory staging becomes `unified_texts_t`, and `stringzilla.cpp` calls `log_environment()` instead of its own copy, which had drifted to printing Goldmont twice.
Scoring counted into one `u32` per needle and then dense-scanned the whole row per document. At the 346,205-needle XLSum vocabulary that is 2.64 MiB of traffic against a 4,898-byte mean document - 557x amplification - and 348 MiB of rows live across 264 blocks against a 50 MiB L2. The right size was never the vocabulary: a document touches a few hundred needles. A block now counts into its own shared table. Under one slot per needle the needle index is the slot, so there is no hash and no probe and the cost matches the dense row it replaces; above that the table is hashed with a per-block overflow row catching what will not seat, and only a document that overflows pays a pass over the vocabulary. Four of the five benchmark slices allocate no overflow row at all. Contributions accumulate as fixed-point integers rather than through a `cub::BlockReduce`. Integer addition is associative, so the block total no longer depends on the order its lanes finish in - a stronger promise than the fixed reduction shape the header used to publish, and one that survives a different grid size. Measured 1 distinct result over 20 chunk-boundary placements at 2.8e-15 relative error, against 3 results at 9.7e-08 for `f32`. Skipping the slots a document never touched also retires a NaN: the dense scan evaluated `substrings_bm25_term` for every needle, so an empty haystack under `length_normalization = 1` computed `0/0` and poisoned the score. On the CPU the ordering that the header publishes is reached by a radix sort over the touched keys instead of an insertion sort quadratic in them, and past half the vocabulary by an ascending walk of the row itself. Scores stay bit-identical, which the split-haystack equality test pins. Entire-vocabulary CUDA scoring goes from 944.9 to 3972.8 MB/s cased and 703.6 to 1567.7 uncased. The counting columns are the control and reproduce the previous table within 1.6x, which is what makes the score columns comparable.
`SZ_HELPER_AUTO` gains a `constexpr` qualifier from C++20 onward, but SVE and RVV vectors are sizeless and therefore non-literal: a `constexpr` function may neither name one in its signature nor hold one in a local until C++23. Every SVE and RVV translation unit consequently failed to compile at exactly the language level the test suite and CI default to, while C++17 and C++23 both built clean. The NEON headers tripped over the same qualifier for a second reason - a `static` local inside a `constexpr` function is itself a C++23 extension, which `-Werror` promotes to a hard failure. Both are a misuse of the macro rather than a property of the targets. Its own comment opens "A scalar helper", and on a vector helper the qualifier is inert in every configuration: below C++23 it fails to compile, and at C++23 it can never fold, since no constant expression of a sizeless type exists. So the fix removes the qualifier where it never did anything, reusing the `SZ_HELPER_INLINE` that was already there, rather than teaching the macro to detect the target. Scoping was measured before landing. Applying `always_inline` to all 627 helper sites cost `test/uncased.cpp` 7.4s -> 243.8s of compile time and 367KB -> 3.96MB of object code; confined to the 101 vector sites it sits inside run-to-run noise and the object is marginally smaller. Validated at C++17, C++20 and C++23 across armv8-a, armv8-a+aes+sha2, armv9-a+sve2+sve2-aes, rv64gcv, rv64gcv_zvkned_zvknhb and x86-64.
Several benchmark rows measured or named something other than what they advertised. `sz_hash_westmere` sat behind `SZ_USE_HASWELL`, so a Westmere-only build lost the hashed-container row entirely; the guard now names both families, since the fastest x86 pairing genuinely mixes a Westmere hash with a Haswell comparator and either can be compiled out alone. Two rows reported `neon` while running the crypto-extension kernels `sz_hash_state_*_neonaes` and `sz_sha256_state_*_neonsha`. The ISA ladders had also drifted behind the headers. `sz_equal` and `sz_order` were missing Westmere, `sz_order` also SVE, the streaming hash and `sz_fill_random` were missing SVE2-AES, multi-seed hashing was missing both V128 tiers, and the whole cipher sweep omitted POWERVSX, RVV-crypto and both V128 tiers - so on those targets the benchmark reported nothing while the kernel existed. `STRINGWARS_TOKENS` rejected the very spellings its own error message told the user to type: the parser accepts `lines` and `words`, the message named `line` and `word`, and two file headers documented the invalid singular as their default. Four more headers documented a default that contradicted the code. Also drops the two UTF-8 wrapper structs stranded in `token.cpp` when those benchmarks moved to `utf8_traverse.cpp`, and `cpu_cycles_per_second`, which had no caller anywhere in the tree.
The suite had no written contract, so a driver's suffix meant whatever its author intended. `test/stringzilla.hpp` now states what each tier costs and may assume, and the drivers were moved to match: `_unit` is fixed-cost known-answer vectors, pinned at every `SZ_TESTS_MULTIPLIER`; `_equivalence` owns randomness and scales; `_safety` drives malformed and boundary inputs; `_all` walks the backend table and holds no assertions of its own. Ten drivers contradicted their own suffix - `test_memory_stability_unit` encoded the violation in its default argument, reading `scale_iterations` from the header. They are renamed or split so the claim holds, and it is now observable rather than asserted: every remaining `_unit` driver reports the same wall time at multiplier 1 and 4, while the renamed `_equivalence` drivers scale. Coverage gaps closed behind that. Sweeps that were written per-ISA by hand had drifted: intersect was missing SVE, uncased ordering SVE2, cased-search RVV, LASX and POWERVSX, and the whole RVV-crypto hash and SHA256 tier was absent. `find`, `hash` and `sort` had no `_safety` driver at all; both new ones initially passed with a planted defect, because they only called dispatched entry points, which resolve to Haswell on x86 and never reach the serial kernel - they now sweep every compiled backend by name. `sz_lookup_init_lower`, `sz_lookup_init_ascii` and the four `sz_find_byte_from` variants had no C++ caller anywhere and are now driven. Comments were brought back in line with the code they describe: driver names that outlived their comments, backends cited by a number the project never assigns, seeds that stopped being replayable once `seed_generator_for_test` mixed in the test name, and 65 trailing comments restating the kernel already spelled on the line. CI ran neither C++11 nor C++14 despite building both, while `find.cpp` gates a block on C++17; the language floor now runs.
The README told readers StringZilla has no Aho-Corasick automaton and pointed them at hyperscan and pyahocorasick instead. It has one - goto-completed, two-tier, running across CPU cores or a CUDA GPU, reachable as `szs.Substrings` and `szs::Substrings` - and its own README describes it in detail. The functionality table omitted the row too. Paths and names had drifted behind two renames. There is no `scripts/` directory, and the test and benchmark files lost their prefixes, so `bench/utf8_iterate.cpp` - a file that never existed under any name - was cited as the source of six benchmark tables, each of which actually belongs to one of `utf8_traverse`, `utf8_scan` or `utf8_segment`. Six more tables named a corpus, `leipzig1M_en.txt`, that is nowhere in the tree. `CONTRIBUTING.md` documented `STRINGWARS_MAX_TOKENS`, which nothing reads, listed 7 of the 13 variables the harness does read, passed a dataset positionally to a harness that declares `argv` unused, and named a `_serial` test target that CMake never defines. The bindings misdescribed their own behaviour. `translate` documented a bytes-to-bytes mapping where the code requires single-character strings, and promised a whole new string where a `start`/`end` pair returns only that slice. Go's `Utf8Count` claimed to match `utf8.RuneCount`, while it counts non-continuation bytes and so returns 0 where `RuneCount` returns 2. Swift documented a `szScope` that does not exist and a `reset()` that restores the original seed, though the seed is never retained. Every claim here was checked against the code, and where the code could answer, by running it: the two `translate` cases and the `Utf8Count` divergence are measured, not read. All 294 doctests still pass and the pytest suite still collects.
Pulls the released `sz_string_reserve` shrink no-op, the `arrow_strings_tape` offset-overflow guards, the forward-iterator `static_assert`, and the CPython helper linkage fix into the multi-pattern branch. The only conflict was the test roster: `main` added `test_string_reserve_unit` where this branch had renamed the two neighbouring drivers to `_equivalence`, so the new unit keeps its name beside the renamed pair.
`SZ_HELPER_AUTO` carries `constexpr` from C++20 onwards, which is what lets a CUDA kernel reach a scalar helper without the C layer naming an execution space. A helper that touches an intrinsic can never be constant-evaluated, though, and a `constexpr` function with no constant-evaluated path is ill-formed: Clang and MSVC reject the definition, GCC 12 too, and only GCC 13+ softens it to `-Winvalid-constexpr`. Carrying the qualifier across every backend cost nine CI jobs at once - both GCC 12 runners, both Windows runners under C3615, StringZillas-CPUs and StringZillas-CUDA whose host compiler is `gcc-12`, and both Clang runners through the Rust `cc` build, which has no warning flags of its own. The 259 intrinsic-reaching helpers move to `SZ_HELPER_INLINE`, so no ISA backend header holds a `SZ_HELPER_AUTO` at all, and `-Wno-invalid-constexpr` leaves the warning set rather than spreading to `build.rs`: nothing that includes these headers needs a flag to compile them. The function-local tables in seven backends become legal where they stand, since only the qualifier made them a C++23 extension. The compiler enumerates the whole set - `-Winvalid-constexpr` under GCC 13+, once per ISA - and its silence is the proof the sweep is complete. The four `substrings` helpers a CUDA walk calls keep the qualifier and now earn it: they take `span<char const>` rather than punning a `span<byte_t const>` into the C iterators, so nothing in the bodies blocks constant evaluation. `szs_substrings_replace_*` spells its offsets `sz_size_t`, the width the engines address the output tape with. Apple Clang had refused the mismatch outright, `sz_u64_t` being a distinct type from `size_t` there rather than merely a same-width one.
`stringzilla_shared` no longer defines `SZ_OVERRIDE_LIBC=1`, so it exports no `memcpy`, `memmove`, `memset`, `memchr`, or `memfrob` of its own. The shim units default that macro to `SZ_AVOID_LIBC`, which this target already set to zero, so deleting the line is the whole change. This changes the behaviour of an installed artifact: `LD_PRELOAD`-ing `libstringzilla.so` no longer accelerates a program's LibC calls. That capability moves to `stringzilla_bare`, which keeps the override and whose whole purpose is replacing LibC. What the default buys is that anything may now link the library - symbol interposition no longer redirects `libstdc++`'s internal `memset` into a dispatch table that is not up yet, which used to fault the process before `main`.
`SZ_DYNAMIC_DISPATCH=1` externs the `sz_*` entry points the engines call - `sz_copy` from the rewriting paths among them - and nothing resolved them for the dynamic variant. ELF defers that to load time, so Linux never complained; Mach-O demands every symbol at link time, and `libstringzillas_cpus_shared.dylib` failed to link once the macOS job got far enough to try. `stringzillas_cpus_shared` now links `stringzilla_shared`, which no longer interposes LibC and is therefore safe to depend on. The archive gets no such line: it resolves nothing itself, and whoever links it brings its own core - tests and benchmarks keep linking `stringzilla_static`, whose `-static` executables cannot take a `.so` at all and whose capability set is fixed at build time rather than by a shared library's own runtime detection. A dynamic StringZillas without a dynamic core is now refused at configure time, where it used to surface as `cannot find -lstringzilla_shared` deep into the build.
ashvardanian
force-pushed
the
main-multipattern
branch
2 times, most recently
from
August 19, 2026 14:55
46bb882 to
0abd25c
Compare
Two causes behind the remaining red jobs, neither of them new to this branch - both were only reachable once the earlier `constexpr` fix let those jobs get further. `SZ_HELPER_AUTO` drops `constexpr` on MSVC. Its bit-scan and byte-swap intrinsics are not constant-evaluable, so `sz_u64_ctz` and everything that reaches one transitively - `sz_size_bit_ceil`, `sz_u64_bits_reverse`, the folded rune iterators, the uncased search - is rejected with C3615, an error no `/wd` silences. Demoting the callers is not the answer: the folded iterators are exactly what a CUDA walk reaches through `--expt-relaxed-constexpr`, and stripping the qualifier there would cost the device its access. MSVC hosts neither `nvcc` nor a fold through those intrinsics, so it loses nothing by seeing the plain inline helper. The substrings fixtures now build haystacks that carry their needles. Ten letters spell the whole corpus, so at `SZ_TESTS_MULTIPLIER=0.05` - what the QEMU jobs run - ten needles drawn against three haystacks matched nothing, and the construction fixture said so through its own guard. The cover and rewriting fixtures share the shape without the guard: every check there sits inside a loop over the match set, so an empty one passes having compared nothing. `random_haystacks_with_needles_` plants one needle per haystack at rotating offsets, which keeps the matches interior, where covers and neighbours mean something.
ashvardanian
force-pushed
the
main-multipattern
branch
from
August 19, 2026 16:25
0abd25c to
f04027c
Compare
`stringzillas_*_shared` reaches `cuGetErrorName` and `cuInit`, which resolve out of `cuda` rather than `cudart`. An ELF shared object carries them undefined until load time, so only a Windows DLL reports it. This is the repair 15b336c already made on the per-target call, lost when that call folded into `define_stringzillas_shared` and only `cudart` came along. `RelWithDebInfo` matched both nvcc flag lists and handed `ptxas` a `-G` beside `-O2`, which it refuses with "Optimized debugging not supported". `-G` stays with the unoptimized `Debug` build and `RelWithDebInfo` takes `-lineinfo`, the same source correlation carried through optimized device code. `MANIFEST.in` drops the two entries whose files left with the old benchmarks and that every sdist warned about. The parallel Python targets declare the NumPy their modules already call `import_array()` against, which the build requirements named and the install requirements did not.
Mark the scalar UTF-8 helpers `constexpr` through `SZ_HELPER_AUTO`, the qualifier `--expt-relaxed-constexpr` needs to resolve a host helper from a kernel at all. Two constructs blocked constant evaluation and both are gone. The `sz_u8_t const *` reinterpretation of the input pointer cannot appear in a constant expression, so `sz_utf8_byte_at_` takes the sign off one byte at a time instead. The 16-entry lead-length table becomes three comparisons, matching the form `sz_rune_decode_unchecked` already uses, so the two can never disagree on how far a lead byte advances and the uncased search path carries no lookup at all. A `__global__` in the substrings suite calls `sz_rune_decode` and `sz_rune_encode` directly, so a helper that stops being reachable stops the file from compiling. Without it `nvcc` resolves the call from inside the folded cursor without a diagnostic, every multi-byte codepoint decodes as a run of malformed single bytes, and an uncased device search matches nothing outside ASCII while every host backend still agrees with the oracle. The device backend is compiled on every CUDA runner and executed on none of them, so a build error is the only form this can be caught in without a GPU.
`pack_hot_children_` advanced `lowest_free_cursor_` after every placement, so its walk was amortized linear as the docblock claimed. `pack_cold_children_` read the same cursor but never moved it, so every cold parent re-scanned the whole occupied prefix from the same low point - linear per state, quadratic per build. Above the 107,520-state hot cap every state is cold, so the largest dictionaries paid all of it. The cold path now publishes the first vacancy its scan found. Every slot below that one is occupied by construction, and the row's own lowest claim is `candidate`, so the cursor steps past it only when the row landed there - no arena is wasted and no transition changes. On 245 MB of XLSum, serial construction over the full 346,205-needle vocabulary drops from 9.23 s to 1.82 s cased and 9.35 s to 1.00 s uncased, and cost per cold state stops climbing across the sweep. The benchmark also stops building the serial automaton twice per cell: the measured row leaves one behind, and only a filtered-out row still owes a build.
The three CUDA domains gave three different answers to where an output must live. Fingerprints refused host memory in one engine and validated nothing in the other - the one the C API picks on its common fast path. Similarities probed the result matrix and staged it when it was host memory. Substrings staged some outputs and required others to be host-writable, because a host loop filled them after the launch. Now one rule holds everywhere: on a GPU scope every buffer an operation reads or writes must be device-accessible, unified or plain device memory, and page-locked host memory is host memory here since the driver reports it as such. `check_device_accessible_memory` and `check_device_accessible_sequence` state that once and every engine calls them. `is_device_accessible_memory(nullptr)` now answers false, which closes the hole that let a null output pass every check. Every staging buffer and drain copy is gone rather than made conditional: the match and output-byte staging, the three BM25 staging vectors, the rewrite offsets, and the similarities host-scatter fallback with its `results_staging_`. Deleting that fallback turns on the direct single-word Myers kernel, which was gated on the results being device-accessible and which no Python caller had ever reached. The rewrite also stops zero-filling the caller's offsets before a capacity refusal, so a refused CUDA call names the size it wanted - which the CPU engines and the C header already promised. Heavy operation bodies move behind a `*_described_` seam so they stop being instantiated once per container shape, and `is_tape_like` gives `corpus_is_ascii_` one `find_byteset` over a whole tape instead of one launch per element.
A GPU scope now refuses host buffers, which left Python users no way to satisfy it for the arrays they supply themselves: `to_device` only accepts a `Strs`, so `score_bm25` was unreachable on a GPU scope without CuPy or Torch. `szs.unified_array(shape, dtype)` closes that - it wraps the unified allocator the engines already use, and the returned NumPy array frees the allocation through its capsule base. Element width comes from `PyDataType_ELSIZE`, since NumPy 2 made `PyArray_Descr::elsize` private. Result matrices and rewrite offsets are allocated device-side on a GPU scope, and an `out=` argument goes through `parse_device_buffer`, which reads `__cuda_array_interface__` and `__dlpack__` and refuses anything else. `sz_device_memory_mismatch_k` raises `BufferError` rather than falling through to a bare `RuntimeError`. `try_swap_to_unified_allocator` now separates the two ways it could fail: a non-`Strs` is a type error and says so, where before every wrong-typed argument on a CUDA build was reported as a device-memory mismatch. `document_lengths` is validated for shape and length against the haystack count, which was a live out-of-bounds read. The dependency groups become one family, `tests-oracles` and `tests-cuda`, both installed by `pip install --group` and both kept out of cibuildwheel's `test-requires`, which runs under QEMU. CuPy is not an oracle but a producer: a source of device buffers that is not our own `unified_array`.
`test_substrings_cuda_memory_safety` covered inputs only, and only `try_count`; nothing tested the fingerprints or similarities guards at all. Each domain now has a memory-safety driver covering unified inputs accepted, host inputs refused, host outputs refused per verb, pure device outputs accepted, and pinned outputs refused. The fingerprints one runs twice, once per CUDA engine, choosing an aligned and a misaligned `dimensions` so the C ABI routes to each. Both new drivers were checked against a deliberately disabled guard first, so they are not vacuous. `substrings_rewrite_tape_t` held `std::string` and `std::vector` while driving the CUDA engine, and the benchmark's rewrite offsets and BM25 weights were host buffers fed to CUDA engines; all are unified now. `device_vector` is `safe_vector` with `device_alloc` rather than a new container, since `std::vector` moves elements through host stores. On the Python side, `test_alignment_out_buffer_matches_returned_matrix` branched on `szs.__capabilities__` - a build fact - to choose a per-scope assertion, and rebound `out_buffer` so its trailing oracle assertion no longer mentioned `out=` at all. It now parametrizes over `DEVICE_NAMES` like its siblings. `device_float32_array` in `szs_helpers` allocates BM25 weights and lengths for whichever scope a test runs on.
`@section <id> <title>` binds the first token as the anchor and the rest as the visible
title, so `@section Device Memory` silently anchored "Device" and titled the section
"Memory" - the heading read wrong and nothing could link to it. Thirteen sections were
written that way. Each now carries a lowercase snake_case id prefixed by its domain, since
ids share one namespace across the project and two headers both document "Buffer Sizing".
The pre-commit hook grows a check for it. A Title-Case first token is the exact signature
of a missing id, and an id with no title renders an empty heading, so both are refused.
The existing decorative-separator rule only read comments, so a banner printed at runtime -
`printf("\n=== Fingerprints ===\n")` - passed straight through it; it now reads string
literals too, requiring the run of dashes or equals to sit against whitespace so a rule used
as a divider is caught while `"---"` as test data is not. The four such banners in the tree
became plain sentences.
The tree is also brought to clang-format-23 across every tracked C/C++ file outside the
submodule, which was 18 files of drift. One docstring in `python/stringzilla/memory.c` had a
two-fragment literal with no trailing `//` to pin it, which the formatter shreds; it is
re-broken so every line carries its own tail.
`pack_cold_children_` scanned interior vacancies until one fit the parent's whole child mask, and the vacancies a packed arena leaves behind are mostly singletons no multi-byte row can ever cover. The scan therefore re-walked the same unfillable holes for every later row, and candidates per row grew linearly with the state count - 10.8 at 0.85M states, 545.2 at 6.2M - making the phase quadratic overall. It was 70-94% of build time and the only reason a multi-million-state dictionary took tens of seconds. The search now gives up after `max_interior_probes_k` rejections and settles the row on the arena frontier, tracked in `claim_slot_` as one past the highest claimed slot. Every slot from there up is unclaimed by construction, so the fallback always fits, and anchoring the row's smallest child byte on the frontier rather than an anchor byte past it leaves nothing stranded behind it. Packing drops from 23.9s to 1.52s on 400K DNA k-mers at 6.2M states, and from 0.58s to 0.19s on 318K English words. The published slot count, the transition-table footprint, and the emitted match lists are all unchanged - byte-identical on both corpora against the previous automaton.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Most Frequent 1% of the Vocabulary
aho_corasick<DFA>aho_corasick<ContiguousNFA>aho_corasickdaachorseregex::find_iterstringzillas::Substrings<1xSPR>stringzillas::Substrings<16xSPR>stringzillas::Substrings<H100>pyahocorasickahocorasick_rsre.finditerstringzillas.Substrings<1xSPR>stringzillas.Substrings<16xSPR>stringzillas.Substrings<H100>Most Frequent 10% of the Vocabulary
aho_corasick<DFA>aho_corasick<ContiguousNFA>aho_corasickdaachorseregex::find_iterstringzillas::Substrings<1xSPR>stringzillas::Substrings<16xSPR>stringzillas::Substrings<H100>pyahocorasickahocorasick_rsstringzillas.Substrings<1xSPR>stringzillas.Substrings<16xSPR>stringzillas.Substrings<H100>Least Frequent 1% of the Vocabulary
aho_corasick<DFA>aho_corasick<ContiguousNFA>aho_corasickdaachorseregex::find_iterstringzillas::Substrings<1xSPR>stringzillas::Substrings<16xSPR>stringzillas::Substrings<H100>pyahocorasickahocorasick_rsstringzillas.Substrings<1xSPR>stringzillas.Substrings<16xSPR>stringzillas.Substrings<H100>Least Frequent 10% of the Vocabulary
aho_corasick<DFA>aho_corasick<ContiguousNFA>aho_corasickdaachorseregex::find_iterstringzillas::Substrings<1xSPR>stringzillas::Substrings<16xSPR>stringzillas::Substrings<H100>pyahocorasickahocorasick_rsstringzillas.Substrings<1xSPR>stringzillas.Substrings<16xSPR>stringzillas.Substrings<H100>Entire Vocabulary
aho_corasick<DFA>aho_corasick<ContiguousNFA>aho_corasickdaachorseregex::find_iterstringzillas::Substrings<1xSPR>stringzillas::Substrings<16xSPR>stringzillas::Substrings<H100>pyahocorasickahocorasick_rsstringzillas.Substrings<1xSPR>stringzillas.Substrings<16xSPR>stringzillas.Substrings<H100>