# USearch for C

> Search, attention and sparse algebra behind one C ABI, in RAM, VRAM or on NVMe, from C

Building the library takes CMake 3.28 and a C++20 compiler, while calling it takes any C99 compiler.
It exposes all four modalities, each served by the engine in RAM, on NVMe and on NVIDIA GPUs; `-D USEARCH_BUILD_STUB=1` adds a brute-force stand-in under the same names.

## Installation

```sh
cmake -B build_release -D CMAKE_BUILD_TYPE=Release -D USEARCH_BUILD_SHARED=1 -D USEARCH_DENSE_DTYPES=11 && cmake --build build_release --config Release
```

That leaves both libraries in `build_release/`, carrying all eleven dense storage types; precompiled ones ship on the [releases page](https://github.com/unum-cloud/USearch/releases).
Put `c/` on the include path, or `include/` once installed, where the C headers sit beside the C++ ones, and link either library.
One header, `usearch/usearch.h`, pulls in every modality: `dense_index.h`, `sequence_index.h`, `shapes_index.h` and `geospatial_index.h`.
`ctest --test-dir build_release -R usearch_test_c` runs the smoke test in `c/test.c`.

`-D USEARCH_INSTALL=1` adds the CMake package and a `.pc` file, and `cpack` then archives the preview SDK.
It holds the stripped shared library, the C headers, the C++ wrappers in `cpp/usearch/`, the package, the `.pc` and the license, but neither the archive nor the header-only core.

```sh
cmake -B build_release -D CMAKE_BUILD_TYPE=Release -D USEARCH_BUILD_SHARED=1 -D USEARCH_INSTALL=1 -D USEARCH_DENSE_DTYPES=11
cmake --build build_release --config Release
cpack --config build_release/CPackConfig.cmake # usearch-<version>-<platform>-sdk.tar.gz, a ZIP on Windows
```

Consumers link it as `usearch::shared` after `find_package(usearch)`, or through `pkg-config --cflags --libs usearch`.
`cmake --install build_release --component static` adds `usearch::static`, and `--component core` the header-only templates as `usearch::header`.

`-D USEARCH_BUILD_STUB=1` adds a brute-force stand-in under the same library names in `build_release/stub/`.
It answers every verb of every modality by exhaustive scan with the C++ standard library alone, so a binding links and runs its suite in seconds without building the engine.
Its search is exact, its traversal statistics are invented, and its checkpoints are its own memory image, readable by nothing else.

## Quickstart

Every verb takes one argument struct, initialized by designator, so every field left out keeps its default.
Every struct opens in the same order: the handle, `executor`, `outcome`, then `completion`; inputs follow, then outputs and the `arena` they land in, then knobs, and last `progress` and `stats`.
Structs carry no size or version field, so they grow source-compatibly only, and the major version changes whenever one does.

```c
#include <stdio.h>           // For printing.
#include <usearch/usearch.h> // For the win 🚀

int main(void) {
    usearch_outcome_t outcome = {0};
    usearch_executor_t executor = NULL;
    usearch_arena_t arena = NULL;
    usearch_dense_index_t index = NULL;
    usearch_key_t keys[2] = {{.bytes[15] = 1}, {.bytes[15] = 2}};
    float vectors[2][128] = {{0}};
    usearch_matches_view_t matches = {0};

    // Threads live in the executor, the data lives in the index, and outputs land in the arena.
    usearch_executor_for_ram(&(usearch_executor_for_ram_t) {.executor = &executor, .outcome = &outcome});
    if (outcome.status != usearch_status_success_k) goto cleanup;
    usearch_arena_init(&(usearch_arena_init_t) {.arena = &arena, .executor = executor, .outcome = &outcome});
    if (outcome.status != usearch_status_success_k) goto cleanup;
    usearch_dense_init_ram(&(usearch_dense_init_ram_t) {.index = &index,
                                                        .outcome = &outcome,
                                                        .metric = usearch_dense_metric_angular_k,
                                                        .quantization = usearch_scalar_bf16_k,
                                                        .dimensions = 128});
    if (outcome.status != usearch_status_success_k) goto cleanup;

    // Populate:
    usearch_dense_insert(
        &(usearch_dense_insert_t) {.index = index,
                                   .executor = executor,
                                   .outcome = &outcome,
                                   .count = 2,
                                   .keys = {keys, sizeof(usearch_key_t)},
                                   .vectors = {vectors, sizeof(vectors[0]), 128, usearch_scalar_f32_k}});
    if (outcome.status != usearch_status_success_k) goto cleanup;

    // Search one query; a zero stride would repeat it across a larger `count`.
    usearch_dense_search(&(usearch_dense_search_t) {.index = index,
                                                    .executor = executor,
                                                    .outcome = &outcome,
                                                    .count = 1,
                                                    .queries = {vectors[0], 0, 128, usearch_scalar_f32_k},
                                                    .matches = &matches,
                                                    .arena = arena,
                                                    .wanted = 10});
    if (outcome.status != usearch_status_success_k) goto cleanup;
    printf("Found %zu matches\n", matches.counts[0]);

cleanup:
    if (outcome.message) fprintf(stderr, "%s: %s\n", usearch_status_name(outcome.status), outcome.message);
    usearch_dense_free(index);
    usearch_arena_free(arena);
    usearch_executor_free(executor);
    return outcome.status == usearch_status_success_k ? 0 : 1;
}
```

Every call here blocks, so its outcome is final by the time it returns.

## Executors

An executor owns the threads or device streams that run calls, the traversal scratch they walk with, and the pinned bounce buffer staging passes through.
One executor serves any number of indexes, of any quantization and metric, since its scratch is sized from geometry, grown to the widest asked of it, and kept between calls.
A VRAM executor's `scratch_limit` caps each stream's run, which every call's walk scratch and outputs carve from, and a call needing more reports `usearch_status_out_of_memory_k`.
One call holds an executor at a time, while a handle admits any number of concurrent callers: the executor serializes, never the index.

```c
int32_t const sockets[2] = {0, 1};
usearch_executor_for_ram(&(usearch_executor_for_ram_t) {.executor = &executor,
                                                        .outcome = &outcome,
                                                        .ram = {sockets, 2}, // NUMA nodes whose cores it spawns on
                                                        .threads = 0});      // zero takes 7/8 of those cores
```

Threads interleave across the named domains, so a half-size executor takes half of every socket rather than all of one, and the default stops at 7/8 because full occupancy starves ordered publication.
`usearch_executor_for_vram` opens one stream per device, or borrows the caller's, and `usearch_executor_for_nvme` spawns one thread per submission ring.
An application with its own ForkUnion pool lends it with `usearch_executor_on_pool` instead, though the engine refuses borrowed pools with `usearch_status_unsupported_k` for now.
An executor's kind must match the residency of every handle it drives, and a mismatch reports `usearch_status_device_k`.

A `NULL` executor is accepted by every verb and runs the call on the caller's own thread, final when it returns whatever its `completion`, which suits size lookups and other short calls.
The exception is a VRAM verb that takes an arena: that arena must be a GPU arena on the index's first device, and it serves only calls on the executor that made it.
A second call on a held executor finishes `usearch_status_contended_k`, so callers wanting real parallelism give each thread its own executor.
`usearch_executor_drain` waits until no call holds the executor, then sets its threads' mood until the next call: `usearch_mood_parked_k`, the default, lets idle threads sleep in the kernel, `usearch_mood_spinning_k` keeps them hot.
A blocking verb called from inside a progress handler or predicate of a call on the same executor reports `usearch_status_contended_k`, and so does a drain on the thread running that call's body; a drain from any other of its threads would wait on itself.
`usearch_executor_free` waits out the call holding the executor before it releases the threads; host arenas the executor made stay valid until freed, and GPU arenas it made must be freed before it.

### Completion & Deferred Calls

Every call reports through its `outcome` and names its `completion` in the same struct, so no verb has a second, asynchronous twin.
`usearch_completion_blocking_k`, the default, returns with the final status stored and every output readable.
`usearch_completion_deferred_k` returns once submitted, as an io_uring entry does, and the call finishes on its executor.
A thread pool, a GPU stream or a remote server can serve the same verbs without changing a call site.

Submission copies the argument struct, but everything it points to stays borrowed until the outcome is final.
One call holds an executor at a time, so calls overlap by running on different executors, each with its own outcome and arena:

```c
usearch_matches_view_t matches[2];
usearch_outcome_t outcomes[2] = {{0}};
for (size_t shard = 0; shard != 2; ++shard)
    usearch_dense_search(&(usearch_dense_search_t) {
        .index = shards[shard],
        .executor = executors[shard],
        .outcome = &outcomes[shard],
        .completion = usearch_completion_deferred_k,
        .count = 32,
        .queries = {queries, 128 * sizeof(float), 128, usearch_scalar_f32_k},
        .matches = &matches[shard],
        .arena = arenas[shard],
        .wanted = 10});
while (usearch_outcome_status(&outcomes[0]) == usearch_status_deferred_k ||
       usearch_outcome_status(&outcomes[1]) == usearch_status_deferred_k)
    do_other_work(); // polling alone finishes both
```

The status reads `usearch_status_deferred_k` until the call is final, even for a call refused at submission.
Then the library stores the message and the final status with release ordering, and touches nothing the struct names after that.
The caller reads it with acquire ordering: `usearch_outcome_status` polls and `usearch_outcome_wait` blocks on an atomic wait, both returning the status.
Those two are the only safe readers of a deferred outcome, since the status word is atomic only by their contract.
Never wait on a call's own outcome from inside its progress handler or predicate, since the call finishes only after they return.
A call ends in one order: its arena and the counts of the handles it names, then the pool, then the executor, then the message and the status.
So a caller that reads success may resubmit on the same executor or free a handle at once.

A call submitted while another holds the executor finishes at once with `usearch_status_contended_k`, and nothing queues, so submitting never allocates.
Every `*_free` waits until each executor call naming its handle has ended, so a handle may be dropped at any moment, though never from a callback of a call naming it.
Waits, drains and frees only wait: none of them runs another thread's call.

### Inclusive and Exclusive Executors

A caller-inclusive pool, a single thread or a WebAssembly build needs its caller to make progress, so a deferred call on it runs at submission and is final when the verb returns.
A caller-exclusive pool, which `usearch_executor_for_ram` and `usearch_executor_for_nvme` spawn whenever they spawn more than one OS thread, delegates a deferred call to its own thread 0.
That thread runs the call in the caller's place and forks the other threads as the caller would, so the call finishes while the caller does something else, one call per executor at a time.
Progress handlers and predicates of a delegated call therefore run on the pool's threads.
VRAM executors run every verb at submission, since each verb reads its device results back on the host before it ends.
Together these let one manager thread in `int main()` overlap a RAM, an NVMe and a VRAM executor: it submits a deferred call to each and polls their outcomes.

## Keys

A `usearch_key_t` is sixteen opaque bytes, enough for a UUID, a content hash, or a 64-bit integer with room to spare, and spelled canonically as a lowercase dashed UUID.
Integer identities sit big-endian in the last bytes with the rest zeroed, which is how the other bindings convert integers:

```c
usearch_key_t key_of(uint64_t value) {
    usearch_key_t key = {{0}};
    for (unsigned byte = 0; byte != 8; ++byte) key.bytes[15 - byte] = (uint8_t)(value >> (byte * 8));
    return key;
}
```

Keys need not be unique: an index admits up to its `values_per_key` entries under one key, unbounded by default, and the count verbs report how many a key holds.
A batch passes its keys as a `usearch_keys_view_t`, a pointer and a byte stride, where a zero stride repeats the first key for every row.

## Inputs and Outputs

### Arenas

Every output whose length scales with the call lands in an arena: the matches of a search, every vector stored under one key, every pair a sparse product scored, a monolith, a cursor's batch.
An arena is memory the library allocates and the caller passes to call after call, each of which rewinds it and writes views into it, growing only when the new outputs outsize everything it held.
An executor makes it where it computes, and a `NULL` executor makes host pages for calls on the calling thread; an arena of host pages may outlive the executor that made it.
Nothing is allocated until a call writes into it, and it keeps its high-water mark, so a steady loop allocates once.
Only fixed-size reports, such as stats, descriptions and `written` counts, stay fields the caller points at.
Arenas keep the interface at one verb per operation, never a twin per output style.

```c
usearch_arena_t arena = NULL;
usearch_arena_init(&(usearch_arena_init_t) {.arena = &arena, .executor = executor, .outcome = &outcome});
for (size_t i = 0; i != 2; ++i) {
    usearch_rows_view_t rows = {0};
    size_t rows_count = 0;
    usearch_dense_get(&(usearch_dense_get_t) {.index = index,
                                              .executor = executor,
                                              .outcome = &outcome,
                                              .key = keys[i],
                                              .vectors = &rows,
                                              .written = &rows_count,
                                              .arena = arena,
                                              .outputs_scalar = usearch_scalar_f32_k});
    if (outcome.status != usearch_status_success_k) break;
    // `rows.data` holds `rows_count` vectors `rows.stride` bytes apart,
    // valid until the next call given `arena`
}
usearch_arena_free(arena);
```

Any verb may reuse any arena of its residency, so a search and a get can share one, as long as the first call's views are not read after the second starts.
To keep one call's outputs alive past the next, pass the next call another arena.
An arena serves one call at a time: a call holds it until just before its final status, and a second call naming it meanwhile finishes `usearch_status_contended_k`.
A `NULL` arena, or inputs pointing into the arena a call is about to rewind, refuse with `usearch_status_invalid_argument_k`, and a call that fails leaves its outputs `NULL`.
A GPU arena grows, fills and frees on the first stream of the executor that made it, so it serves only calls on that executor, and any other executor, `NULL` included, is refused with `usearch_status_device_k`.
Free a GPU arena before the executor that made it.

### Rows & Scalars

A batch of rows is a `usearch_rows_view_t`: a base pointer, a byte stride, the values per row, and the scalar kind they hold, converted on the way in where it differs from storage.
A zero stride repeats the first row, as in BLAS, and the call's `count` alone is the batch axis.
Vectors are quantized on insertion into the index's `quantization`, from `f64` down to `i8`, `u8`, the four 8- and 6-bit floats `e5m2`, `e4m3`, `e3m2` and `e2m3`, and packed 4- and 1-bit integers.
The library compiles the first `USEARCH_DENSE_DTYPES` storage types of `bf16`, `u1x8`, `i8`, `e4m3`, `f16`, `u8`, `e5m2`, `i4x2`, `u4x2`, `e3m2` and `e2m3`, and refuses the rest with `usearch_status_unsupported_k`.
The precompiled binaries carry all eleven, while a CMake build takes the first four unless configured with `-D USEARCH_DENSE_DTYPES=11`.

USearch comes pre-packaged with NumKong, bringing SIMD-accelerated distance kernels for x86, Arm, RISC-V and GPUs.
The dense metric kinds are:

- `usearch_dense_metric_inner_k` - Inner Product distance, $1 - \sum_i a_i b_i$, unnormalized and scale-sensitive.
- `usearch_dense_metric_angular_k` - Cosine distance, $1 - \frac{\sum_i a_i b_i}{\sqrt{\sum_i a_i^2} \sqrt{\sum_i b_i^2}}$, in $[0, 2]$ and blind to magnitude.
- `usearch_dense_metric_euclidean_k` - Euclidean distance, $\sqrt{\sum_i (a_i - b_i)^2}$, zero only for identical vectors.
- `usearch_dense_metric_hamming_k` - Hamming distance, $\sum_i a_i \oplus b_i$, over `usearch_scalar_u1x8_k` bit sets.
- `usearch_dense_metric_jaccard_k` - Jaccard distance, $\frac{\sum_i a_i \oplus b_i}{\sum_i a_i \lor b_i}$, in $[0, 1]$, over `usearch_scalar_u1x8_k` bit sets.

`usearch_scalar_name`, `usearch_dense_metric_name` and their `_from_name` inverses spell these sets the way every binding parses them.

### Search

A search writes a `usearch_matches_view_t` into its arena: `columns` packed columns for each of its `count` queries, query `i` holding `counts[i]` matches from `keys[i * columns]` and `scores[i * columns]` on, nearest first.
`wanted` bounds the columns, the "k" in "kANN".
The graph's `connectivity` is frozen when the index is created, while the `expansion_factor`, the breadth of the candidates a walk explores, is chosen per call.
Higher expansion values buy recall at the cost of latency; zero derives a default from the connectivity.

```c
usearch_traversal_stats_t stats = {0};
usearch_dense_search(&(usearch_dense_search_t) {.index = index,
                                                .executor = executor,
                                                .outcome = &outcome,
                                                .count = 1,
                                                .queries = {query, 0, 128, usearch_scalar_f32_k},
                                                .matches = &matches,
                                                .arena = arena,
                                                .wanted = 10,
                                                .expansion_factor = 128,
                                                .stats = &stats});
```

`.traversal = usearch_traversal_scan_k` scores every published entry instead of walking, the exact ground truth a walk approximates, on RAM, NVMe and NVIDIA VRAM.
AMD and Apple VRAM handles refuse scans with `usearch_status_unsupported_k`.
Every verb that walks the graph fills the optional `stats` with the members it visited, the distances it computed, and its restarts, a connectivity health signal rather than a cost.

## Metrics

### Dense Indexes

A dense index holds fixed-width vectors under one metric, with the `dimensions`, `quantization`, `connectivity` and `values_per_key` its constructor froze.

## Residency

An index lives in RAM, VRAM or on NVMe, and every lifecycle verb names its action and its residency instead of taking a tag a residency might not support:

- `init` starts empty and mutable: `usearch_dense_init_ram`, `usearch_dense_init_vram`.
- `load` copies a checkpoint in: `usearch_dense_load_ram`, `usearch_dense_load_vram`, and their `_buffer` forms.
- `view` serves one in place, read-only: `usearch_dense_view_ram`, `usearch_dense_view_nvme`.

The options name where an index lives, never who walks it.
`usearch_ram_options_t` names NUMA nodes, one shard per node, so a dual-socket box gets two graphs whose walks stay local; payload blocks claim huge pages where a pool is provisioned.
`usearch_vram_options_t` names device ordinals, one shard per device and never a replica, so an index outgrowing them fails rather than paging to the host.
Both take a `shard_capacity`, the entries one shard may hold, zero meaning the 32-bit ID space: an insertion past it opens another generation of shards on the same nodes or devices, and only the newest generation takes insertions.
`usearch_nvme_options_t` names paths: a filesystem path selects a kernel I/O engine, a transport identifier like `trtype:PCIe traddr:0000:5e:00.0` selects SPDK in a `-D USEARCH_WITH_SPDK=1` build, and a FUSE mount over a bucket reads as an ordinary path.

Memory follows residency: an index keeps its payload where its lifecycle verb put it, an executor keeps its scratch next to its threads or devices, and an arena lives where the executor that created it computes.
So outputs come back in RAM from a RAM executor and in VRAM from a VRAM executor, and moving them elsewhere is the caller's own copy.
Callers lend memory, never a way to allocate it: no allocator callback crosses the interface, because large systems hold too many kinds of memory for one ABI to describe them all.

Every dense verb is exported for every residency, and the engine answers what a residency lacks with `usearch_status_unsupported_k`:

- RAM answers every verb but `usearch_dense_view_ram` and `usearch_dense_search_stream_init`, and exact scans serve `multiply` and `attend` too.
- VRAM answers insertion, refinement, absorbing, counting, `get`, search, `multiply`, `maxsim`, attention and the serialization family, `save_buffer` included, while cursors answer `usearch_status_unsupported_k`, as [GPUs](#gpus) explains.
  NVIDIA devices answer exact scans, while AMD and Apple ones refuse them with `usearch_status_unsupported_k`.
  Apple's Metal answers attention `usearch_status_unsupported_k` too, since its kernels cannot call the caller's value lookup.
  A sequence, shape or geospatial VRAM index answers `save_buffer` with `usearch_status_unsupported_k` for now, and those modalities load from memory into RAM alone, through `load_ram_buffer`.
- NVMe answers every read verb as RAM does, `sample` and `health` included, refuses every mutating one with `usearch_status_read_only_k`, and refuses `save` and `save_buffer`, since a view already is its checkpoint.

### GPUs

A VRAM index lives on the devices its `vram` options name, one shard per device, and a VRAM executor drives it through one stream per device.
Passing your own streams lets USearch work join the rest of your pipeline, though the engine refuses borrowed streams with `usearch_status_unsupported_k` for now:

```c
#include <cuda_runtime.h>

cudaStream_t stream;
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
int32_t const device = 0;
usearch_executor_t gpu = NULL;
usearch_executor_for_vram(&(usearch_executor_for_vram_t) {
    .executor = &gpu, .outcome = &outcome, .vram = {&device, 1}, .streams = (void* const*)&stream});

usearch_dense_index_t index = NULL;
usearch_dense_init_vram(&(usearch_dense_init_vram_t) {.index = &index,
                                                      .outcome = &outcome,
                                                      .metric = usearch_dense_metric_inner_k,
                                                      .quantization = usearch_scalar_bf16_k,
                                                      .dimensions = 1024,
                                                      .vram = {&device, 1}});
usearch_dense_reserve(&(usearch_dense_reserve_t) {
    .index = index, .executor = gpu, .outcome = &outcome, .capacity = 1000000});
```

The library classifies every pointer it is given through the driver, so inputs may live anywhere:
rows already on the device are read in place, pinned host rows upload directly, and only pageable host rows pass through the executor's pinned bounce buffer.
Pin the rows you upload often, and the bounce disappears:

```c
float* rows = NULL;
cudaHostAlloc((void**)&rows, 4096 * 1024 * sizeof(float), cudaHostAllocDefault);
// ... fill 4096 rows ...
usearch_dense_insert(&(usearch_dense_insert_t) {.index = index,
                                                .executor = gpu,
                                                .outcome = &outcome,
                                                .count = 4096,
                                                .keys = {row_keys, sizeof(usearch_key_t)},
                                                .vectors = {rows, 1024 * sizeof(float), 1024, usearch_scalar_f32_k}});
```

An arena lives where the executor that created it computes, so on a VRAM executor every output an arena holds is __device memory__.
Feed it straight into your next kernel, or bring it back yourself, into the layout you want.
Here the packed rows `get` returns land in a pinned host table with a wider row pitch, in one strided copy:

```c
usearch_rows_view_t rows_on_device = {0};
size_t rows_count = 0;
usearch_arena_t arena = NULL;
usearch_arena_init(&(usearch_arena_init_t) {.arena = &arena, .executor = gpu, .outcome = &outcome});
usearch_dense_get(&(usearch_dense_get_t) {.index = index,
                                          .executor = gpu,
                                          .outcome = &outcome,
                                          .key = row_keys[0],
                                          .vectors = &rows_on_device,
                                          .written = &rows_count,
                                          .arena = arena,
                                          .outputs_scalar = usearch_scalar_f32_k});

size_t const table_pitch = 1280 * sizeof(float);
float* table = NULL;
cudaHostAlloc((void**)&table, rows_count* table_pitch, cudaHostAllocDefault);
cudaMemcpy2DAsync(table, table_pitch, rows_on_device.data, rows_on_device.stride, rows_on_device.stride, rows_count,
                  cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
usearch_arena_free(arena);
```

A search's matches land in its arena too, so on a VRAM executor they are device memory: every shard ranks its own entries and the first device merges those rankings into the arena, so on NVIDIA and AMD no match crosses through the host, while on Apple the host merges them in the unified memory both share.
A device that runs out of memory for an arena reports `usearch_status_out_of_memory_k` rather than falling back to the host, so where an output lands never depends on how full a device is.
Free a device arena once your own device work on its outputs has finished, and before the executor that made it.

## Serialization

Checkpoints ride the SafeGraphs interchange format, and the path's extension picks the layout: a `.safegraphs` path is one monolith of every shard, a `.safegraphs.json` path is a manifest beside one monolith per shard.
Loading reads the content, never the extension, so either form loads into any shard count, part `p` loading into shard `p` modulo the shard count.

```c
usearch_bytes_view_t path = {(uint8_t const*)"index.safegraphs", 16};
usearch_dense_save(&(usearch_dense_save_t) {.index = index, .executor = executor, .outcome = &outcome, .path = path});

usearch_dense_index_t copy = NULL, mapped = NULL;
usearch_dense_load_ram(&(usearch_dense_load_ram_t) {
    .index = &copy, .executor = executor, .outcome = &outcome, .path = path});
usearch_dense_view_ram(&(usearch_dense_view_ram_t) {.index = &mapped, .outcome = &outcome, .path = path});
```

A loaded index is a mutable copy, while a viewed one serves the file in place through a memory mapping and refuses every mutating verb with `usearch_status_read_only_k`.
The engine's RAM shards cannot serve a mapping yet, so it answers `usearch_dense_view_ram` with `usearch_status_unsupported_k`, while `usearch_dense_view_nvme` serves checkpoints in place today.
A cancelled or failed save leaves a file without its header or parts without a manifest, which every reader rejects and the caller removes, and a cancelled or failed load creates no handle.
A checkpoint's geometry and metric can be read without loading it:

```c
usearch_dense_metadata_t metadata = {0};
usearch_dense_metadata_read(&(usearch_dense_metadata_read_t) {
    .executor = executor, .outcome = &outcome, .path = path, .metadata = &metadata});
```

An index can also be saved into memory as one monolith of every shard, which loads into any shard count.
The bytes land in an arena, so the call needs no size estimate upfront:

```c
usearch_bytes_view_t monolith = {0};
usearch_dense_save_buffer(&(usearch_dense_save_buffer_t) {
    .index = index, .executor = executor, .outcome = &outcome, .monolith = &monolith, .arena = arena});
usearch_dense_load_ram_buffer(&(usearch_dense_load_ram_buffer_t) {
    .index = &copy, .executor = executor, .outcome = &outcome, .buffer = monolith});
```

## Filtered Search

Sometimes you may want to cross-reference search results against some external database or filter them by some criteria.
In most engines, you'd have to manually perform paging requests, successively filtering the results.
In USearch you can pass a predicate to the search, which is applied during graph traversal.

```c
void admits_tenant(usearch_key_t const* keys, usearch_admission_t* admissions, size_t count, void* state) {
    for (size_t i = 0; i != count; ++i)
        admissions[i] = keys[i].bytes[0] == *(uint8_t const*)state ? usearch_admit_k : usearch_reject_k;
}

uint8_t tenant = 7;
usearch_dense_search(&(usearch_dense_search_t) {.index = index,
                                                .executor = executor,
                                                .outcome = &outcome,
                                                .count = 1,
                                                .queries = {query, 0, 128, usearch_scalar_f32_k},
                                                .matches = &matches,
                                                .arena = arena,
                                                .wanted = 10,
                                                .predicate = {admits_tenant, &tenant}});
```

The predicate judges a batch of candidates per call, a quad of them at a time during a walk, so a handler backed by an external store answers each batch in one round trip.

Rejected entries still route the walk, so recall holds at low selectivity, and the predicate judges every scored candidate, concurrently on the executor's threads, so it must be fast and tolerate that.
An exact scan asks it only about entries closer than the worst match a query already holds, which makes the scan the fallback for predicates too selective for a walk to cross.
A host predicate cannot narrow a device walk, so a VRAM search with one reports `usearch_status_unsupported_k`.

Long calls take a `usearch_progress_t`: a handler, which is also the only way to stop one, its state, and how many reports it hears:

```c
usearch_continuation_t stop_when_asked(size_t completed, size_t total, void* state) {
    return *(int const*)state ? usearch_cancel_k : usearch_continue_k;
}

int stop = 0;
usearch_dense_search(&(usearch_dense_search_t) {
    .index = index, .executor = executor, .outcome = &outcome, .count = 1000000, .queries = queries,
    .wanted = 10, .arena = arena, .matches = &matches, .progress = {stop_when_asked, &stop, 20}});
```

Counts are in whatever units the verb consumes, entries for an insertion, queries for a search, shards for a checkpoint, and the handler is never called concurrently or after the call is final.
A call reports at most `steps` times, the last at completion, and zero steps derive to ten.
A batch verb hands its engine one slice per step, and a call without a handler runs as one slice, since every engine bounds its own scratch.
A stopped call reports `usearch_status_cancelled_k` and leaves its outputs `NULL`.
An insertion claims every entry before its first report, so its handler only reports, and all its entries publish at once.

## Refinement

`usearch_dense_refine` re-selects the neighbors of every stored entry against the finished graph, improving those chosen while the graph around them was incomplete.
A cheap insertion pass followed by one refinement reaches a better graph in less total time than one expensive pass, so a wide `expansion_factor` belongs here rather than in `usearch_dense_insert`.
Re-selection is idempotent and claims no identity, so a cancelled pass leaves the graph correct and merely less improved, still re-links any entry no search would reach, and answers `usearch_status_cancelled_k`.
The index keeps nothing of a stopped pass, so a caller looping until the graph settles calls again:

```c
size_t changed = 0;
do {
    usearch_dense_refine(&(usearch_dense_refine_t) {.index = index,
                                                    .executor = executor,
                                                    .outcome = &outcome,
                                                    .changed = &changed,
                                                    .expansion_factor = 256});
} while (outcome.status == usearch_status_success_k && changed > 100);
```

Searches may run throughout on a RAM handle, while a VRAM handle rewrites neighbors in place and admits no overlapping search.
Insertions never overlap a refinement: one running refuses it with `usearch_status_contended_k`, one arriving meanwhile is refused the same way and lands nothing, and so on VRAM is a save, which holds the handle the same way for its own duration.
Nothing waits, so a service schedules refinements between its insertion bursts and retries a refused call.

A refinement that must continue where a cancel stopped is a handle the caller holds instead.
`usearch_dense_refinement_init` opens it, optionally absorbing donors first, and the index refuses insertions and saves until it is finished.
Each `usearch_dense_refinement_next` continues it on the executor it was opened with, holding that executor only for the call, and answers `usearch_status_cancelled_k` while more is left.
It reports the running totals, the nodes re-selected so far, those whose neighbors changed, and the telemetry, so a caller stops calling once the changed share stops paying for another call.
`usearch_dense_refinement_finish` re-links what a stopped pass leaves, reports, reopens the index and frees the handle.

### Absorbing & Splitting

`usearch_dense_absorb` moves an array of donor handles into one index, one donor after another, keeping their graphs rather than re-inserting their vectors.
Donors are consumed and nulled as they go, so the array itself says which handles the caller still owns and must free, and a cancel stops only the closing refinement, answering `usearch_status_cancelled_k`.
Each donor's shards append to the index's as appended parts with their edges rebased, then a refinement of every shard that took one walks each of its nodes across into the other part from landmarks drawn there and re-selects it, so a group costs one such refinement per donor, and its heal re-selects from the landmarks every node no earlier one links to.
A donor must share the index's residency, scalar, class, metric and geometry, else the call answers `usearch_status_config_mismatch_k`, and an NVMe view answers `usearch_status_read_only_k`.
An insertion running on the index or on the next donor refuses the call with `usearch_status_contended_k`, leaving that donor and the rest live, and an insertion into the index arriving meanwhile is refused the same way, while nothing may search a donor or start an insertion into it.
A VRAM-built donor joins a RAM index through a save and a load.

`usearch_dense_split` is the inverse, emptying one index into `count` receivers, empty handles of its residency, scalar, class, metric, geometry and shard count.
It grows one graph-local part per receiver from roots drawn from `seed`, breadth-first over the links and capped at an equal share, lands part `p` of shard `s` in shard `s` of receiver `p` with edges rebased and edges across parts dropped, and refines the nodes that lost an edge.
`usearch_dense_refinement_init` takes `receivers` in place of `donors` for the resumable form, never both, refusing the pair with `usearch_status_invalid_argument_k`.
RAM and NVIDIA and AMD VRAM indexes split, Apple and NVMe ones answer `usearch_status_unsupported_k` and `usearch_status_read_only_k`, and the stand-in partitions by brute force.

## Linear Algebra

### Sparse Multiplication

`usearch_dense_multiply` is an approximate $A \times B^T$ in the shape of BLAS `?GEMM`, with the indexed rows as $B$, reporting every pair the walks scored rather than only the nearest.
The output pattern is whatever the walk touched, so its cost is roughly `count × expansion × connectivity × dimensions`, sublinear in the indexed entries, and the pairs it never reaches are the low and negative dots.
Only the walks learn how many pairs there are, so keys, products and Apache Arrow offsets land in the arena as one ragged output:

```c
usearch_matches_tape_view_t products = {0};
size_t required = 0;
usearch_dense_multiply(&(usearch_dense_multiply_t) {.index = index,
                                                    .executor = executor,
                                                    .outcome = &outcome,
                                                    .count = m,
                                                    .left = {left_rows, 128 * sizeof(float), 128, usearch_scalar_f32_k},
                                                    .products = &products,
                                                    .required = &required,
                                                    .arena = arena,
                                                    .wanted = 256});
for (size_t row = 0; row != m; ++row)
    for (size_t pair = products.offsets[row]; pair != products.offsets[row + 1]; ++pair)
        accumulate(row, products.keys[pair], products.scores[pair]);
```

Pairs arrive in traversal order without duplicates, a non-zero `wanted` keeps each left row's first pairs, and `required` reports how many there were before the cut.
The linear-algebra verbs run on `usearch_dense_metric_inner_k` alone, since products and softmaxes need logits rather than distances, and any other kind refuses with `usearch_status_config_mismatch_k`.

### Late Interaction

`usearch_dense_maxsim` is ColBERT scoring, and it is what `values_per_key` above one is for: a document is one key holding its token rows, and its score sums each query row's best match among them.
The graph stays over individual tokens, so no set-to-set metric is evaluated during the walk, and documents need not share a row count.
Queries share one `rows_per_query`, or `query_offsets` gives each its own run of rows:

```c
usearch_matches_tape_view_t documents = {0};
size_t written = 0;
usearch_dense_maxsim(&(usearch_dense_maxsim_t) {
    .index = index,
    .executor = executor,
    .outcome = &outcome,
    .count = 4,
    .queries = {query_tokens, 128 * sizeof(float), 128, usearch_scalar_f32_k},
    .rows_per_query = 32,
    .matches = &documents,
    .written = &written,
    .arena = arena,
    .wanted = 10});
```

Query `i` spans `documents.offsets[i]` up to `documents.offsets[i+1]`, largest first, and `written` counts every query's documents on the host, even when the arena lives in VRAM.
The walks only nominate keys, and every entry of a nominated key is scored, so a walk may miss a key but never scores one short, and widening the expansion narrows what it misses.

### Attention

`usearch_dense_attend` folds Flash Attention's online softmax over the keys each walk scores, with the graph replacing the sweep over key tiles and no top-k cut, so the tail the walk scores brings the result toward exact attention.
One index holds one layer's one key-value head, row `t * head_count + h` is head `h` of token `t`, and the `head_count` query heads of grouped-query attention share one walk per token.
The index stores no values, so `value_lookup` resolves the keys of every scored batch to rows as wide as the keys, in the index's own scalar kind:

```c
void value_rows_of(void* state, usearch_key_t const* keys, size_t count, usearch_scalar_t scalar, size_t columns,
                   void const** rows) {
    for (size_t i = 0; i != count; ++i) rows[i] = find_value_row(state, keys[i]); // a NULL row refuses the call
}

usearch_dense_softmax_state_t carry[64 * 8];
usearch_rows_view_t attended = {0};
usearch_dense_attend(&(usearch_dense_attend_t) {.index = index,
                                                .executor = executor,
                                                .outcome = &outcome,
                                                .count = 64 * 8,
                                                .queries = {query_rows, 128 * 2, 128, usearch_scalar_bf16_k},
                                                .head_count = 8,
                                                .value_lookup = value_rows_of,
                                                .value_lookup_state = &values,
                                                .carry = carry,
                                                .outputs = &attended,
                                                .arena = arena});
```

On RAM `value_lookup` is a host function, and on VRAM a `__device__` one whose address comes from `cudaMemcpyFromSymbol` on NVIDIA or `hipMemcpyFromSymbol` on AMD; passing the wrong kind is undefined, never refused.
`carry` holds each row's running maximum and partition, which merges this branch with the recent window kept out of the index, attention sinks, or a second index of one context:
$m = \max(m_1, m_2)$, $a_i = l_i e^{m_i - m}$, $l = a_1 + a_2$, $o = (a_1 o_1 + a_2 o_2) / l$.
Attention never writes the graph, so it overlaps insertions freely, the publication watermark acting as the causal mask.

## Retrieval and Streaming

`usearch_dense_get` writes every vector stored under one key into an arena, shard by shard and in insertion order within each, as the [Arenas](#arenas) example shows; an absent key leaves `vectors->data` `NULL`.
A cursor walks the graph, every entry in storage order or breadth-first from `starts`, its keys and, when asked, each entry's `links` copied into an arena batch by batch:

```c
usearch_dense_cursor_t cursor = NULL;
usearch_dense_cursor_init(&(usearch_dense_cursor_init_t) {.index = index, .outcome = &outcome, .cursor = &cursor});
usearch_keys_view_t batch = {0};
for (size_t written = 1; written != 0 && outcome.status == usearch_status_success_k;)
    usearch_dense_cursor_next(&(usearch_dense_cursor_next_t) {.index = index,
                                                              .executor = executor,
                                                              .outcome = &outcome,
                                                              .cursor = cursor,
                                                              .keys = &batch,
                                                              .written = &written,
                                                              .arena = arena,
                                                              .wanted = 1024});
usearch_dense_cursor_free(cursor);
```

Each shard's publication watermark is read as the walk enters it, so an entry once passed is never revisited, while a later shard may carry entries published after the walk began.
Opened at `starts`, a cursor yields their own entries first, in the order given, then breadth-first the entries they link to, each entry once, so the first entries pulled carry the starts' immediate links.
`links` is a `usearch_keys_tape_view_t`, entry `i` spanning `offsets[i]` up to `offsets[i+1]`, and the vectors themselves come from `usearch_dense_get`.

When a predicate is selective enough that no `wanted` can be guessed upfront, a search stream keeps pulling matches in distance order:

```c
usearch_dense_search_stream_t stream = NULL;
usearch_dense_search_stream_init(&(usearch_dense_search_stream_init_t) {
    .stream = &stream, .index = index, .outcome = &outcome, .query = {query, 0, 128, usearch_scalar_f32_k}});
do {
    usearch_dense_search_stream_next(&(usearch_dense_search_stream_next_t) {.index = index,
                                                                            .executor = executor,
                                                                            .outcome = &outcome,
                                                                            .stream = stream,
                                                                            .matches = &matches,
                                                                            .arena = arena,
                                                                            .wanted = 10});
    // ... keep the matches that pass, and stop once enough did ...
} while (outcome.status == usearch_status_success_k && matches.counts[0] != 0);
usearch_dense_search_stream_free(stream);
```

No residency implements streams yet, so `usearch_dense_search_stream_init` answers `usearch_status_unsupported_k` today.
A stream keeps the walk's state between pulls, never a thread, so any number of them may stay open on one executor.
Cursors and streams keep no hold on the index either: every pull passes the index it was opened on, any other is refused with `usearch_status_invalid_argument_k`, and an overlapping pull reports `usearch_status_contended_k`.

## Introspection

One call reports an index's frozen geometry and its current counters:

```c
usearch_dense_index_description_t description = {0};
usearch_dense_describe(&(usearch_dense_describe_t) {
    .index = index, .executor = executor, .outcome = &outcome, .description = &description});
printf("%zu of %zu entries, %zu bytes, %zu dimensions, class: %s\n", description.size, description.capacity,
       description.memory_usage, description.metadata.dimensions, usearch_cpu_class_name(description.cpu_class));
```

`usearch_dense_count` reports how many entries a key holds, as when a long document is split into chunks.
`usearch_dense_health` audits how much of the graph a search reaches, sweeping forward from as many landmarks as one search draws:

```c
usearch_reachability_t reachability = {0};
usearch_dense_health(&(usearch_dense_health_t) {
    .index = index, .executor = executor, .outcome = &outcome, .reachability = &reachability});
// `reachability.reachable` of `reachability.audited` entries are found by some search
```

`usearch_capabilities_compiled` and `usearch_capabilities_available` name the SIMD and GPU capabilities the library was built with and the machine offers, and `description.cpu_class` or `description.gpu_class` names the class an index dispatched to, which `usearch_cpu_class_name` and `usearch_gpu_class_name` spell.
`usearch_version` reports the linked library's version, which may differ from the `USEARCH_VERSION_MAJOR` family of macros the caller compiled against.

## Errors

The library never logs or aborts: every failure is a status in the call's `usearch_outcome_t`, with a static message carrying the detail the status cannot.
Messages are never freed, and are `NULL` exactly on success.
`usearch_status_name` spells a status as its lowercase dashed name, like `invalid-argument`, and `usearch_status_from_name` reads one back.
The statuses worth branching on:

- `usearch_status_unsupported_k`: the verb is declared but this residency or build does not implement it, so no retry can help.
- `usearch_status_read_only_k`: a mutating verb on a viewed handle.
- `usearch_status_contended_k`: refused rather than waited for, like an executor or an arena serving another call, an insertion meeting a refinement, absorption or clear of its index, or a second VRAM insertion beside the first.
- `usearch_status_cancelled_k`: a progress handler stopped the call, which published nothing.
- `usearch_status_device_k`: an executor, arena or handle on another residency or device than the call needs, or a GPU arena named with another executor than the one that made it.
- `usearch_status_invalid_argument_k`: a required pointer left `NULL`, inputs inside the arena the call rewinds, a name no table spells.

The set describes the project's failures, not only this library's exits, so a binding validating before the call reports what the library would have.
We aim for no `NULL` that switches behaviour, no bitmask that mixes unrelated options, and no field that repeats what another already says.

## Platform Notes

### Sequence Indexes

A sequence index searches byte strings, UTF-8 text, DNA, RNA and proteins under edit distances and alignments, in RAM only.
Variable-length entries travel as Apache Arrow tapes, `usearch_tape_view_t`: the elements end to end, and `count + 1` offsets around them, 32-bit or 64-bit as its `layout` says, so sequences are bytes with byte offsets:

```c
usearch_sequence_index_t sequences = NULL;
usearch_sequence_init_ram(&(usearch_sequence_init_ram_t) {
    .index = &sequences, .outcome = &outcome, .metric = usearch_sequence_metric_levenshtein_utf8_k});

char const tape[] = "kittensittingmitten";
uint64_t const offsets[4] = {0, 6, 13, 19};
usearch_key_t words[3] = {{.bytes[15] = 1}, {.bytes[15] = 2}, {.bytes[15] = 3}};
usearch_sequence_insert(&(usearch_sequence_insert_t) {.index = sequences,
                                                      .executor = executor,
                                                      .outcome = &outcome,
                                                      .count = 3,
                                                      .keys = {words, sizeof(usearch_key_t)},
                                                      .sequences = {.data = tape,
                                                                    .layout = usearch_tape_u64_k,
                                                                    .offsets = {.u64 = offsets},
                                                                    .scalar = usearch_scalar_u8_k}});

uint64_t const query_offsets[2] = {0, 6};
usearch_sequence_search(&(usearch_sequence_search_t) {.index = sequences,
                                                      .executor = executor,
                                                      .outcome = &outcome,
                                                      .count = 1,
                                                      .queries = {.data = "sitten",
                                                                  .layout = usearch_tape_u64_k,
                                                                  .offsets = {.u64 = query_offsets},
                                                                  .scalar = usearch_scalar_u8_k},
                                                      .matches = &matches,
                                                      .arena = arena,
                                                      .wanted = 3});
```

The kinds are Levenshtein over bytes or UTF-8 code points, Needleman-Wunsch with a linear gap, and overlap, the share of fixed-width windows one side misses of the other.
Needleman-Wunsch takes `usearch_sequence_weights_t`, a 256×256 substitution matrix with a gap cost, and the overlap kind takes `window_widths` in bytes, each refused by every other kind with `usearch_status_config_mismatch_k`.
The verbs mirror the dense ones: `init_ram`, `init_vram`, `load_ram`, `load_vram`, `view_ram`, `reserve`, `clear`, `insert`, `refine`, `search` with a predicate, `get` writing tapes into an arena, `cursor_next` writing keys and their links, `absorb`, `split`, `save`, `metadata_read`, the `_buffer` forms of `load_ram`, `save` and `metadata_read`, and `describe`, while `count` also reports a key's total and largest byte lengths.
A key already holding `values_per_key` entries rejects the whole insertion with `usearch_status_overflow_k`.
The engine serves every kind, and answers `view_ram` with `usearch_status_unsupported_k`; `refine`, `absorb` and `split` serve RAM indexes, while `sample`, `health` and cursors with `starts` and `links` serve RAM indexes and NVMe views.
The stand-in serves every kind by textbook dynamic programming, quadratic and unbanded, and reports alignment similarities negated.

### Shape Indexes

A shape index searches 3D point clouds of any size, like molecules, conformers and protein backbones, in RAM or VRAM.
Its tape is points end to end, `x`, `y` and `z` interleaved, with offsets that count points rather than scalars:

```c
usearch_shapes_index_t conformers = NULL;
usearch_shapes_init_ram(&(usearch_shapes_init_ram_t) {.index = &conformers,
                                                      .outcome = &outcome,
                                                      .metric = usearch_shapes_metric_wasserstein_procrustes_k,
                                                      .quantization = usearch_scalar_f16_k});

float const atoms[5][3] = {{0, 0, 0}, {1.1f, 0, 0}, {0, 0, 0}, {1.5f, 0, 0}, {0, 1.5f, 0}};
uint64_t const point_offsets[3] = {0, 2, 5}; // a 2-atom and a 3-atom entry
usearch_shapes_insert(&(usearch_shapes_insert_t) {.index = conformers,
                                                  .executor = executor,
                                                  .outcome = &outcome,
                                                  .count = 2,
                                                  .keys = {molecules, sizeof(usearch_key_t)},
                                                  .shapes = {.data = atoms,
                                                             .layout = usearch_tape_u64_k,
                                                             .offsets = {.u64 = point_offsets},
                                                             .scalar = usearch_scalar_f32_k}});
```

Chamfer and Hausdorff pair every point with its nearest partner after aligning the principal axes under all eight sign flips, keeping the best.
Wasserstein-Procrustes pairs points by an exact transport plan, alternated with the best rotation or reflection.
Gromov-Wasserstein compares each side's internal distances.
All four are invariant to rotation, reflection, translation and point order.
A non-zero `points` freezes the count, rejecting insertions of any other with `usearch_status_shape_mismatch_k`.
The verbs are the sequence ones, where `get` converts to an `outputs_scalar` and returns each shape in its inserted pose, and `count` reports a key's total and largest point counts.
The engine serves all four kinds in RAM and on NVMe over `f16` or `f32` points, refusing other quantizations and `view_ram` with `usearch_status_unsupported_k`, and Chamfer and Hausdorff on NVIDIA GPUs; `sample`, `health` and cursors with `starts` and `links` serve RAM indexes and NVMe views.
The stand-in scores Chamfer and Hausdorff in the frame the points arrive in, refuses the two transport kinds with `usearch_status_unknown_kind_k`, refuses quantizations other than `f16` and `f32` with `usearch_status_unsupported_k`, and refuses every VRAM verb with `usearch_status_device_k`.

### Geospatial Indexes

A geospatial index searches observations on and above the Earth, like photos, video frames, drone footage and sensor fixes, in RAM, on NVMe or on NVIDIA GPUs.
Each entry is one `usearch_geospatial_observation_t` of eleven doubles in natural units: latitude and longitude in WGS-84 degrees, altitude in metres above the ellipsoid, the horizontal accuracy, the heading, pitch, heading accuracy and field of view of the device's view cone, how far it saw, and the first and last Unix second it covers.
A device that recorded nothing for a field leaves it NaN, never zero, and each NaN only widens the tests that read it: an unknown accuracy counts as exact, an unknown view as looking everywhere, an unknown time as always.
A key may hold several observations, up to the `values_per_key` its creation froze, like a camera's frames or a vehicle's fixes:

```c
#include <math.h> // `NAN` - a field no device recorded

usearch_geospatial_index_t places = NULL;
usearch_geospatial_init_ram(&(usearch_geospatial_init_ram_t) {.index = &places, .outcome = &outcome});

usearch_geospatial_observation_t const stops[3] = {
    // latitude, longitude, altitude, accuracy, heading, pitch, heading accuracy, field, range, begin, end
    {51.5308, -0.1238, NAN, 5, NAN, NAN, NAN, NAN, NAN, NAN, NAN},
    {51.5320, -0.1233, NAN, 5, NAN, NAN, NAN, NAN, NAN, NAN, NAN},
    {51.5033, -0.1196, NAN, 5, NAN, NAN, NAN, NAN, NAN, NAN, NAN},
};
usearch_geospatial_insert(&(usearch_geospatial_insert_t) {.index = places,
                                                          .executor = executor,
                                                          .outcome = &outcome,
                                                          .count = 3,
                                                          .keys = {stop_keys, sizeof(usearch_key_t)},
                                                          .observations = stops});
```

`usearch_geospatial_within` reports every entry a `relation` admits inside the `tolerance`, however many there turn out to be, as one tape per query in the arena:

```c
usearch_geospatial_observation_t const station = {51.5308, -0.1238, NAN, NAN, NAN, NAN, NAN, NAN, NAN, NAN, NAN};
usearch_matches_tape_view_t nearby = {0}; // `nearby.offsets[1]` entries, their chords in metres
usearch_geospatial_within(&(usearch_geospatial_within_t) {.index = places,
                                                          .executor = executor,
                                                          .outcome = &outcome,
                                                          .count = 1,
                                                          .queries = &station,
                                                          .relation = usearch_geospatial_relation_near_k,
                                                          .tolerance = {.metres = 500},
                                                          .matches = &nearby,
                                                          .arena = arena});
```

`usearch_geospatial_search` ranks the `wanted` nearest entries of each query the relation admits instead, where a zero radius ranks by distance alone.
Every score is the chord between the two positions in metres, altitude included, which ranks as the great-circle arc does.
`near_k` admits entries within the radius widened by three standard deviations of both positions' accuracies combined.
`visible_k` also needs a clear line of sight over the handle's `terrain`, and a handle without one refuses it.
`sees_k` needs the entry's view cone to hold the query, which cameras see this place, and `seen_k` the query's cone to hold the entry, what this camera sees; both clear the terrain too when the handle has one.
The tolerance's `seconds` widen the query's interval on both ends, and a query without a time matches entries of any time.
The terrain is a north-up grid of heights above the ellipsoid, borrowed in place or filled on demand by a callback; a checkpoint never carries it, so every load and view names it again, and a GPU index takes borrowed heights alone.
The remaining verbs are the shapes ones, where `get` returns observations within the stored rounding, centimetres and milliseconds among them, `count` reports a key's entries, and `sample` places its draws at their Earth-fixed positions.
The engine serves every relation in RAM, on NVMe and on NVIDIA GPUs, answering `view_ram` with `usearch_status_unsupported_k`; `sample`, `health` and cursors with `starts` and `links` serve RAM indexes and NVMe views.
The stand-in scans every entry exactly, keeps the doubles inserted, refuses every terrain and so `visible_k`, and refuses every VRAM verb with `usearch_status_device_k`.

### What This Binding Does Not Expose

- __Batched admission.__
  The predicate is one call per key; the span-admission form exists only in the C++ layer, since a C callback per candidate batch would cost more than it amortizes.
- __Custom metrics.__
  A metric of your own, over payloads of your own, belongs to the C++ baseline tier, whose index takes any batched metric you write; see [`include/usearch/basic/README.md`](https://github.com/unum-cloud/USearch/blob/main/include/usearch/basic/README.md).

### Design

- __This is the C ABI itself.__
  Every name is the one `c/usearch/` declares: a verb `usearch_dense_<verb>` takes a `usearch_dense_<verb>_t`, and closed sets are C enums with name tables, never strings.
- __The completion chooses, the verb stays.__
  A `NULL` executor runs a verb on the calling thread, an executor with the default completion blocks on it, and `usearch_completion_deferred_k` submits it, all with the same argument struct.
- __Deferred calls report through the same outcome.__
  `usearch_outcome_status` polls and `usearch_outcome_wait` blocks, both with the acquire ordering the final store pairs with.
- __No life of its own.__
  Threads live only in executors, and a deferred call only on its executor's thread 0; every `*_free` waits for the calls naming its handle first.
- __Arenas are lent and reused, never shared.__
  Every output whose length scales with the call lands in an arena the caller passes; each call rewinds it, so the previous call's views go stale the moment the next one starts.
- __Nothing costs anything until it is used.__
  Submitting a call allocates nothing, and an arena allocates only when a call writes into it.
- __One shape for every call.__
  Each verb takes one struct set by designator, so a call names only what it changes and the struct grows without breaking source.
- __Failures are statuses.__
  The library never logs or aborts, and a dense verb a residency lacks answers `usearch_status_unsupported_k` rather than missing from the build.
