C · Linear Algebra

Linear Algebra

#include <usearch/dense_index.h> 3.0.0-rc.1 · 4c3e5023

§ Sparse Multiplication

usearch_dense_multiply is an approximate A×BTA \times B^T in the shape of BLAS ?GEMM, with the indexed rows as BB, 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:

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:

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:

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⁡(m1,m2)m = \max(m_1, m_2), ai=liemi−ma_i = l_i e^{m_i - m}, l=a1+a2l = a_1 + a_2, o=(a1o1+a2o2)/lo = (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.

API

Types

usearch_dense_softmax_state_tThe online-softmax carry for one query row - Flash Attention's running maximum and partition, which is what lets one call's output fold into another's.
dense_index.h:2098 ↗
struct usearch_dense_softmax_state_t

The two are one concept: a maximum without its partition is a number nothing can rescale, so they share an element rather than travelling as two arrays a caller could supply half of.

Fields

maximum_scorefloatLargest logit seen so far - the recurrence's m.
partitionfloatSum of exponentials taken against that maximum - the recurrence's l.
usearch_value_lookup_tResolves count keys to their value rows, one row per key, columns values of scalar each; a NULL row refuses the call.
dense_index.h:2112 ↗
typedef void(*) usearch_value_lookup_t(void *state, usearch_key_t const *keys, size_t count, usearch_scalar_t scalar, size_t columns, void const **rows);

A host function for a RAM handle; for a VRAM one, a __device__ function whose address comes from cudaMemcpyFromSymbol on NVIDIA or hipMemcpyFromSymbol on AMD, its state and rows 16-byte aligned in memory that device reads. No driver query tells the two kinds apart, so passing the wrong one is undefined, never refused.

Functions

usearch_dense_multiplyApproximate sparse product of a run of rows against the indexed vectors - the linear algebra sibling of usearch_dense_search, reporting every pair the traversal scored instead of only the nearest.
dense_index.h:2089 ↗
void usearch_dense_multiply(usearch_dense_multiply_t *arguments);

Approximate Products & Attention

A walk scores many times the pairs a search reports, and each of them is an exact product the caller would otherwise recompute - this verb hands them all back. Pairs arrive in traversal order rather than sorted, and carry no duplicates: the traversal's visited set admits each indexed vector at most once per left-hand row. A wider expansion_factor widens the emitted pattern, that breadth being what decides how much the walk touches.

How many pairs a walk scores is known only once it finishes, so the call writes the keys, the scores and their Apache Arrow offsets in arena as one ragged output, in one call. A non-zero wanted keeps the first pairs each left-hand row's walks scored, shard by shard in placement order, and required says how many there were.

Cost is not O(m · n · k) but roughly O(count × expansion × connectivity × dimensions), sublinear in the indexed vectors, which is the whole point of reading a product off a graph. The result is approximate, and the pairs the walk never reaches are the low and negative dots.

Runs on the inner-product kinds, whose traversal scores are the products themselves. A graph frozen under a distance kind scored a different quantity, so the products it would report are not the ones asked for, and the call refuses with usearch_status_config_mismatch_k.

// BLAS: B is an operand, C is dense m × n, and every pair is computed.
cblas_sgemm(CblasRowMajor, CblasNoTrans, CblasTrans, m, n, k,
            1.0f, a, lda, b, ldb, 0.0f, c, ldc);

// Here: B was built once and is the index, C is ragged and holds only the pairs the walk
// scored, and the executor carries what a cuBLAS handle plus its stream would.
usearch_arena_t arena = NULL;
usearch_matches_tape_view_t c;
usearch_outcome_t outcome = {0};
usearch_arena_init(&(usearch_arena_init_t){.arena = &arena, .executor = executor, .outcome = &outcome});
usearch_dense_multiply(&(usearch_dense_multiply_t){
    .index = index, .executor = executor, .outcome = &outcome,
    .count = m, .left = {a, a_stride, k, usearch_scalar_f32_k}, .products = &c, .arena = arena});
for (size_t row = 0; row != m; ++row)
    for (size_t pair = c.offsets[row]; pair != c.offsets[row + 1]; ++pair)
        accumulate(row, c.keys[pair], c.scores[pair]);

// The next product reuses the arena, and frees nothing until the last one.
usearch_arena_free(arena);

What the comparison leaves out is the whole of BLAS' tuning surface: no transpose flags, no alpha and beta, no leading-dimension triple, and nothing accumulating into C.

Parameters

indexusearch_dense_index_tThe handle supplying the right-hand rows.
executorusearch_executor_tThe executor driving the batch; its kind must match the handle's residency.
outcomeusearch_outcome_t *Output for the outcome, final once its status leaves usearch_status_deferred_k.
completionusearch_completion_tBlocking, the default, or deferred.
countsize_tLeft-hand rows - the batch axis.
leftusearch_rows_view_tThe left-hand rows, multiplied against every indexed vector the walk reaches.
productsusearch_matches_tape_view_t *Output for the pairs inside arena: each one's indexed key and product, query i spanning products->offsets[i] up to products->offsets[i+1].
writtensize_t *The optional output for the pairs written across every query.
requiredsize_t *The optional output for the pairs scored across every query, before wanted cut any.
arenausearch_arena_tMemory for products, rewound by the call.
wantedsize_tThe optional bound on pairs kept per query, the first ones its walk scored; zero keeps every pair.
expansion_factorsize_tThe optional breadth of the candidates, which widens the emitted pattern; zero derives.
traversalusearch_traversal_tHow entries are reached: a walk, the default; no residency multiplies by an exact scan yet.
predicateusearch_predicate_tThe optional predicate, asked about a scored quad of keys at a time; a pair it rejects is neither written nor counted, while the walk still routes through it.
progressusearch_progress_tThe optional handler, called with queries resolved.
statsusearch_traversal_stats_t *The optional traversal telemetry; NULL skips reporting.

See also

BLAS, whose ?GEMM is the dense counterpart this is shaped after: https://netlib.org/blas

Statuses

usearch_dense_attendSoftmax attention of a run of query heads over every indexed key their graph walks score, folded as Flash Attention folds key tiles.
dense_index.h:2247 ↗
void usearch_dense_attend(usearch_dense_attend_t *arguments);

Approximate Products & Attention

One walk per token runs by a search's own frontier and stopping rule, ranking a key by the best of the token's heads, so every head's own region is explored. Every key the walk scores folds into each head's running maximum, partition and weighted value sum, rescaled whenever the maximum rises - Flash Attention's online softmax, with the graph replacing the sweep over key tiles. No top-k cut is taken: the long tail the walk scores is what brings the result toward exact dense attention, and outputs comes back divided by the partition.

Rows of queries, outputs and carry are tasks: row t · head_count + h is head h of token t. One index holds one layer's one key-value head, so the head_count query heads of grouped-query attention share its keys and one walk per token; one head per token is plain multi-head attention.

The index stores no values, so only keys a walk scores are ever resolved to value rows. A RAM handle calls value_lookup on the host once per scored batch; a GPU handle calls it on the device inside the walk. A value row is exactly as wide as the keys and stored in the index's own scalar kind, which every call states as columns and scalar. A NULL row refuses the call with usearch_status_shape_mismatch_k rather than skipping the key: a missing row is a wrong softmax, not a smaller one.

Keep the recent window out of the index, inserting a token once it leaves the window, so the dense window and the graph never share a key. carry holds the softmax's m and l, which is what merges this branch with another over the same queries - the recent window, attention sinks, or a second index of one context. The two travel in one element because a maximum without its partition is a number nothing can rescale. Two branches over the same queries merge as:

m = max(m₁, m₂), aᵢ = lᵢ · e^(mᵢ − m), l = a₁ + a₂, o = (a₁ · o₁ + a₂ · o₂) / l

Runs on the inner-product kinds, since a softmax needs logits and only those kinds score one. A graph frozen under a distance kind refuses with usearch_status_config_mismatch_k.

Never writes to the graph, so any number of these may overlap with concurrent insertions - which is what streaming a growing context through this verb amounts to, the publication watermark acting as the causal mask. Concurrency & Thread Safety

// One layer, one key-value head shared by 8 query heads: row t · 8 + h is head h of token t.
usearch_rows_view_t attended;
usearch_outcome_t outcome = {0};
usearch_dense_attend(&(usearch_dense_attend_t){
    .index = index, .executor = executor, .outcome = &outcome, .count = tokens * 8,
    .queries = {query_rows, dimensions * 2, dimensions, usearch_scalar_bf16_k}, .head_count = 8,
    .value_lookup = value_rows_of, .value_lookup_state = value_cache,
    .outputs = &attended, .carry = carry, .arena = arena});

// The recent window never entered the index, so it merges in by the same rule, counting no key twice.
merge_recent_window(attended.data, attended.stride, carry, tokens * 8, dimensions);

Parameters

indexusearch_dense_index_tThe handle holding one layer's keys for one key-value head.
executorusearch_executor_tThe executor driving the batch; its kind must match the handle's residency.
outcomeusearch_outcome_t *Output for the outcome, final once its status leaves usearch_status_deferred_k.
completionusearch_completion_tBlocking, the default, or deferred.
countsize_tQuery rows - the batch axis; row t · head_count + h is head h of token t.
queriesusearch_rows_view_tThe query rows.
head_countsize_tThe optional query heads sharing the index's key-value head, and with it one walk per token; it must divide count, and zero means one.
value_lookupusearch_value_lookup_tResolves the keys of every scored batch to their value rows, on the host or the device.
value_lookup_statevoid *Opaque pointer forwarded into every value_lookup call.
carryusearch_dense_softmax_state_t *The optional online-softmax state, one per query row, lent by the caller; NULL skips it.
outputsusearch_rows_view_t *Output for the attended rows inside arena, count of them, as wide as the keys.
arenausearch_arena_tMemory for outputs, rewound by the call.
outputs_scalarusearch_scalar_tThe optional scalar type the attended rows land as; zero takes the queries' own.
expansion_factorsize_tThe optional breadth of the candidates; zero derives.
score_scalefloatThe optional logit scale, conventionally the inverse square root of the head dimensionality; zero derives it from the index's frozen geometry.
traversalusearch_traversal_tHow entries are reached: a walk, the default; no residency attends by an exact scan yet.
predicateusearch_predicate_tThe optional predicate, asked about a scored quad of keys at a time; a key it rejects folds into no head and resolves no value row, while the walk still routes through it.
progressusearch_progress_tThe optional handler, called with queries resolved.
statsusearch_traversal_stats_t *The optional traversal telemetry; NULL skips reporting.

See also

Flash Attention, whose recurrence this reuses: https://arxiv.org/abs/2205.14135

Statuses

usearch_dense_maxsimLate-interaction scoring of multi-row queries against multi-row entries: every query row takes its best match within a key, and those maxima sum to that key's score.
dense_index.h:2366 ↗
void usearch_dense_maxsim(usearch_dense_maxsim_t *arguments);

Approximate Products & Attention

This is what values_per_key above one exists for. A document is one key holding its token rows, and the walk stays over individual tokens under a plain inner product - no set-to-set metric is ever evaluated, since the grouping happens in the reduction rather than in the graph. That is also why an entry's rows need no ordering and two documents need not share a row count.

Exact for every key it scores: the walks only nominate keys, each query row taking as many tokens as the expansion factor, and every row of a nominated key then folds into that key's per-row maxima. A walk may miss a key none of its rows ranked; a scan nominates every key. Widening the expansion factor narrows what a walk misses.

Runs on the inner-product kinds, since a maximum over logits is what late interaction sums. A graph frozen under a distance kind refuses with usearch_status_config_mismatch_k.

// 4 queries of 32 padded token rows each, scored against documents of any length.
usearch_matches_tape_view_t documents;
usearch_outcome_t outcome = {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, .arena = arena, .wanted = 10});

Parameters

indexusearch_dense_index_tThe handle holding one token row per entry, several sharing a key.
executorusearch_executor_tThe executor to fan out over; its kind must match the residency.
outcomeusearch_outcome_t *Output for the outcome, final once its status leaves usearch_status_deferred_k.
completionusearch_completion_tBlocking, the default, or deferred.
countsize_tQueries - the batch axis; matches spans count of them.
queriesusearch_rows_view_tThe token rows of every query laid end to end, count times rows_per_query of them, or as many as query_offsets spans.
query_offsetssize_t const *The optional offsets of every query's rows, count + 1 of them; NULL keeps rows_per_query.
rows_per_querysize_tRows in every query; uniform, so ColBERT's padding applies as is; unread when query_offsets is set.
matchesusearch_matches_tape_view_t *Output for each query's best-scoring keys, largest first, as many as it found up to wanted, or every key it scored when wanted is zero, packed into arena: query i spans matches->offsets[i] up to matches->offsets[i+1], scores carrying each key's summed maxima.
writtensize_t *The optional output for the keys written across every query, read on the host even when matches lives in device memory.
requiredsize_t *The optional output for the keys scored across every query, before wanted cut any.
arenausearch_arena_tMemory for matches, rewound by the call.
wantedsize_tThe optional bound on keys per query - the "k" over documents, not token rows; zero keeps every key the query scored.
expansion_factorsize_tThe optional breadth of the candidates; zero derives.
traversalusearch_traversal_tHow candidates are nominated: a walk, the default, or a scan nominating every key.
predicateusearch_predicate_tThe optional predicate over keys, asked before any key's rows are reduced.
hintsusearch_matches_view_t const *The optional keys every row's walk of a query starts from ahead of its random landmarks, one row of counts[query] keys per query, as a search writes its matches and takes them, resolved once per query; scores are not read, a scan ignores them, and keys no entry holds are skipped and counted in stats once per query.
progressusearch_progress_tThe optional handler, called with queries resolved.
statsusearch_traversal_stats_t *The optional traversal telemetry; NULL skips reporting.

See also

ColBERT, which introduced this scoring: https://arxiv.org/abs/2004.12832

Statuses