[Rust][C API] cuvsCagraBuild blocks calling thread for full build duration — add stream-based cuvsCagraBuildAsync

[zbennett10]

Environment

• cuVS: 26.2 (crates.io + C library)
• CUDA: 12.4
• GPU: T4 (sm_75, 16 GB HBM)
• Rust: 1.82 (MSRV)

Problem

cuvsCagraBuild — and its Rust wrapper Index::build() — synchronizes the
CUDA stream before returning. The call blocks the calling thread for the full
duration of graph construction.

use std::time::Instant;

let t = Instant::now();
let _index = Index::build(&res, &params, &dataset)?;
println!("{:?}", t.elapsed());

// Measured on T4, 512K × 1024-D float32:
// 1.823s
//
// The thread is blocked. No other work on this CUDA context can proceed.

In a thread-per-GPU architecture — where one OS thread owns the CUDA context
for the lifetime of the process — this means search requests queue up and
receive no response for the entire build duration.

CAGRA's build kernels (NN-descent graph construction) and search kernels (beam
traversal) operate on entirely different data structures. The graph under
construction and the graph being searched are independent.

Root Cause

cuvsCagraBuild calls resource::sync_stream(res) at the end of
raft::neighbors::cagra::build(). This is correct semantics for the synchronous
case — callers can safely read the index after the call returns — but it removes
any possibility of enqueueing the build without blocking, even when the caller
intends to synchronize manually via a CUDA event or cudaStreamSynchronize.

(If cagra::build() also contains any cudaDeviceSynchronize() calls, the
fix is more involved — please confirm. From inspection of the public RAFT source,
it appears to be stream-synchronous only.)

Proposed Fix

C API (new function, additive)

/**
 * @brief Enqueue CAGRA index build on a caller-provided CUDA stream.
 *
 * Returns immediately after submitting all graph construction kernels.
 * The index is not safe to use for search until the caller synchronizes
 * build_stream (e.g., cudaStreamSynchronize(build_stream) or a CUDA event).
 *
 * Concurrent CUDA work on other streams is unaffected.
 *
 * @param[in]  res          cuVS resources (memory pools, handles)
 * @param[in]  params       Index build parameters
 * @param[in]  dataset      Input dataset (host or device via DLPack)
 * @param[out] index        Output index (valid only after stream sync)
 * @param[in]  build_stream CUDA stream for build kernels
 */
cuvsError_t cuvsCagraBuildAsync(
    cuvsResources_t res,
    cuvsCagraIndexParams_t params,
    DLManagedTensor* dataset,
    cuvsCagraIndex_t index,
    cudaStream_t build_stream
);

Rust binding (new API, additive — build() is unchanged)

/// Handle to an in-progress CAGRA index build.
pub struct AsyncBuildHandle {
    index: Index,
    build_stream: ffi::cudaStream_t,
}

impl AsyncBuildHandle {
    /// Synchronize the build stream and return the completed index.
    pub fn wait(self) -> Result<Index>;
}

impl Index {
    /// Enqueue a CAGRA build on `build_stream` and return immediately.
    pub fn build_async<T: Into<ManagedTensor>>(
        res: &Resources,
        params: &IndexParams,
        dataset: T,
        build_stream: ffi::cudaStream_t,
    ) -> Result<AsyncBuildHandle>;
}

Usage pattern

let build_stream: ffi::cudaStream_t = unsafe { create_cuda_stream()? };

// Enqueue build — returns in microseconds
let handle = Index::build_async(&res, &params, &dataset, build_stream)?;

// Searches on an existing index continue on the default stream concurrently.
// When the new index is needed:
let new_index = handle.wait()?;

Backward Compatibility

Index::build() is unchanged. cuvsCagraBuildAsync is a new C symbol —
purely additive to the ABI. No existing call sites break.

Impact

Without an async build API, applications that need continuous search
availability must maintain multiple pre-built index copies to mask rebuild
latency, multiplying HBM usage by the number of copies. Stream-level
concurrency eliminates this trade-off entirely.

See also PR #1839, which addressed the related issue of Index::search()
consuming the index and forcing pool-based workarounds.