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 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:

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

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:

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.