Skip to content

CAGRA_CUDA_OXIDE_FULL_MIGRATION_PLAN.md

Latest commit

 

History

History
2101 lines (1665 loc) · 65.2 KB

CAGRA_CUDA_OXIDE_FULL_MIGRATION_PLAN.md

File metadata and controls

2101 lines (1665 loc) · 65.2 KB
Raw

CAGRA cuda-oxide Full Migration Implementation Plan

For agentic workers: REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (- [ ]) syntax for tracking.

Goal: Make the production CAGRA CUDA backend use cuda-oxide for build, search, delta search, runtime probing, and resident device caches, with the old .cu/cudarc stack kept only as a temporary parity oracle.

Architecture: Promote src/apps/cagra/gpu/cuda_oxide from an experimental smoke path to the implementation behind the existing cagra-cuda backend contract. Move host-only validation, route, memory-budget, and cache-selection helpers into a CUDA-common module so cuda-oxide and the temporary legacy oracle do not depend on each other. Port base graph search and NN-descent build to #[cuda_module] Rust kernels, prove them with local A4000 cargo oxide run smokes, cut service routing over, then remove legacy production references.

Tech Stack: Rust 2021 on nightly nightly-2026-04-03, cuda-oxide cuda-device/cuda-host/cuda-core rev b0774f66, existing Morpheus CAGRA modules, optional temporary cudarc legacy feature for parity, local NVIDIA RTX A4000 CUDA verification.

Execution Context

Run implementation work in the isolated worktree:

cd "/home/shisoft/Code/OSS Projects/cagra-engine-worktrees/Morpheus-cuda-oxide"
git status --short

The plan document lives in the main repository root:

ls "/home/shisoft/Code/OSS Projects/Morpheus/CAGRA_CUDA_OXIDE_FULL_MIGRATION_PLAN.md"

Preserve unrelated dirty files. Do not run git reset --hard, git checkout --, or broad cleanup commands. Use the old cudarc path only behind a legacy feature until parity is proven.

Original State At Plan Start

Already converted and smoke-tested:

  • src/apps/cagra/gpu/cuda_oxide/delta.rs: delta brute-force L2 cuda-oxide kernel.
  • src/apps/cagra/gpu/cuda_oxide/build.rs: exact-L2 base build cuda-oxide kernel.
  • src/bin/cagra_cuda_oxide_smoke.rs: delta smoke.
  • src/bin/cagra_cuda_oxide_build_smoke.rs: exact-build smoke.

Old cudarc/.cu production code at plan start:

  • src/apps/cagra/gpu/cuda/runtime.rs: cudarc context and NVRTC module cache.
  • src/apps/cagra/gpu/cuda/resident.rs: cudarc resident base segment cache.
  • src/apps/cagra/gpu/cuda/search.rs: production base graph search wrapper.
  • src/apps/cagra/gpu/cuda/kernels/search_l2.cu: production base graph search kernel.
  • src/apps/cagra/gpu/cuda/build.rs: production exact and NN-descent build wrappers.
  • src/apps/cagra/gpu/cuda/kernels/build_nndescent_l2.cu: NN-descent build kernel.
  • src/apps/cagra/service.rs: routes CAGRA CUDA build/search/delta calls through gpu::cuda::*.

Execution Status

The isolated implementation branch completed this migration through 51d2c51 refactor(cagra): remove legacy CUDA production path. Historical checkboxes for Tasks 1-5 were normalized after verifying the branch commit history; Task 10 records the fresh local A4000 verification evidence, and Task 11 records the final route/device-residency review.

Completion Gates

The migration is complete only when all gates pass:

  • cagra-cuda enables cuda-oxide CAGRA production code.
  • CAGRA service CUDA build, base search, and delta search call cuda-oxide implementation paths.
  • Base search traversal, visited set, frontier, row states, and top-k run on device.
  • Exact build and NN-descent build use cuda-oxide kernels.
  • Resident base cache uses cuda-oxide DeviceBuffer and honors cache capacity.
  • Delta and base cache budget enforcement remain combined.
  • Temporary legacy code is not reachable from production CAGRA routes.
  • Local A4000 smokes pass with cargo oxide run.
  • Non-CUDA builds still compile and CPU/Vulkan scaffolding is not regressed.

File Responsibility Map

Create:

  • src/apps/cagra/gpu/cuda_common.rs: host-only CAGRA CUDA validation, build selection, memory estimates, descriptor keys, byte helpers, and LRU eviction selection.
  • src/apps/cagra/gpu/cuda_oxide/runtime.rs: cuda-oxide context probe/cache and cache clearing.
  • src/apps/cagra/gpu/cuda_oxide/resident.rs: cuda-oxide resident segment upload and LRU cache.
  • src/apps/cagra/gpu/cuda_oxide/search.rs: cuda-oxide base graph search kernel and launch wrapper.
  • src/bin/cagra_cuda_oxide_search_smoke.rs: exact build plus base graph search smoke.
  • src/bin/cagra_cuda_oxide_nndescent_smoke.rs: NN-descent build smoke.
  • src/bin/cagra_cuda_oxide_e2e_smoke.rs: service-level CAGRA CUDA smoke after route cutover.

Modify:

  • Cargo.toml: make cagra-cuda use cuda-oxide deps, add temporary cagra-cuda-legacy, add smoke binaries.
  • src/apps/cagra/gpu/mod.rs: export production cuda as cuda-oxide and export temporary cuda_legacy only under cagra-cuda-legacy.
  • src/apps/cagra/gpu/cuda_oxide/mod.rs: export runtime, resident, search, build, delta.
  • src/apps/cagra/gpu/cuda_oxide/build.rs: import shared host helpers from cuda_common, add NN-descent.
  • src/apps/cagra/gpu/cuda_oxide/delta.rs: use shared host helpers where names overlap.
  • src/apps/cagra/gpu/cuda/build.rs: temporary legacy import from cuda_common.
  • src/apps/cagra/gpu/cuda/resident.rs: temporary legacy import from cuda_common.
  • src/apps/cagra/gpu/cuda/search.rs: temporary legacy import from cuda_common.
  • src/apps/cagra/service.rs: route through production gpu::cuda::*, which becomes cuda-oxide.
  • src/apps/cagra/tests.rs: update CUDA route assertions to accept cuda-oxide route reasons.
  • src/apps/cagra/gpu/cuda/tests.rs: either move host-only tests to cuda_common tests or gate as legacy.

Remove after parity:

  • src/apps/cagra/gpu/cuda/kernels/*.cu
  • production use of src/apps/cagra/gpu/cuda/runtime.rs
  • production use of src/apps/cagra/gpu/cuda/resident.rs
  • production use of src/apps/cagra/gpu/cuda/search.rs
  • production use of old cudarc build/delta wrappers

Task 1: Establish Baseline And Protect Existing Proofs

Files:

  • Modify: no source files in this task.

  • Test: existing cuda-oxide smoke binaries.

  • Step 1: Capture current worktree status

Run:

cd "/home/shisoft/Code/OSS Projects/cagra-engine-worktrees/Morpheus-cuda-oxide"
git status --short

Expected: dirty files may be present. Do not revert unrelated changes.

  • Step 2: Re-run existing exact-build smoke

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_build_smoke --features cagra-cuda-oxide

Expected output contains:

PASSED: cuda-oxide CAGRA exact build smoke matched CPU exact-L2 oracle
  • Step 3: Re-run existing delta smoke

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_smoke --features cagra-cuda-oxide

Expected output contains:

PASSED: cuda-oxide CAGRA delta smoke matched CPU top-k
  • Step 4: Commit baseline only if this task introduced status files

Run:

git diff --stat

Expected: this task introduces no source changes. Skip commit if no files changed.

Task 2: Split Host CUDA Helpers Into cuda_common

Files:

  • Create: src/apps/cagra/gpu/cuda_common.rs

  • Modify: src/apps/cagra/gpu/mod.rs

  • Modify: src/apps/cagra/gpu/cuda/build.rs

  • Modify: src/apps/cagra/gpu/cuda/resident.rs

  • Modify: src/apps/cagra/gpu/cuda/search.rs

  • Modify: src/apps/cagra/gpu/cuda_oxide/build.rs

  • Test: src/apps/cagra/gpu/cuda_common.rs inline tests.

  • Step 1: Write failing common-helper compile target

Create src/apps/cagra/gpu/cuda_common.rs with the moved host-only API names first:

use std::{
    collections::{HashMap, HashSet},
    hash::Hash,
    mem::size_of,
};

use dovahkiin::types::Id;
use neb::index::vector::CagraBuildAlgo;

use crate::apps::cagra::types::{CagraError, CagraSearchConfig, CagraSegmentDescriptor};

pub const CAGRA_CUDA_BUILD_MAX_GRAPH_DEGREE: usize = 256;
pub const CAGRA_CUDA_BUILD_AUTO_EXACT_ROW_THRESHOLD: usize = 16_384;
pub const CAGRA_CUDA_SEARCH_BLOCK_THREADS: u32 = 256;
pub const CAGRA_CUDA_EMPTY_OUTPUT_ROW: u32 = u32::MAX;
pub const CAGRA_CUDA_DEFAULT_DYNAMIC_SHARED_BUDGET_BYTES: usize = 48 * 1024;

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum CudaCagraBuildAlgorithm {
    ExactL2,
    NnDescentL2,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub struct CagraCudaBuildShape {
    pub row_count: u32,
    pub dimension: u32,
    pub graph_degree: u32,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
pub struct CagraCudaSegmentKey {
    pub device_ordinal: u32,
    pub schema_id: u32,
    pub field_id: u64,
    pub segment_id: u64,
    pub generation: u64,
    pub checksum: [u8; 32],
}

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub struct CagraCudaSearchLaunchResources {
    pub output_count: usize,
    pub hash_table_slots: usize,
    pub shared_mem_bytes: u32,
    pub requires_hash_resets: bool,
}
  • Step 2: Move existing build helpers into cuda_common

Move these exact functions from src/apps/cagra/gpu/cuda/build.rs into cuda_common.rs and keep their signatures:

pub fn select_cuda_build_algorithm(
    row_count: usize,
    dimension: usize,
    graph_degree: usize,
    requested: CagraBuildAlgo,
) -> CudaCagraBuildAlgorithm

pub fn validate_cuda_build_inputs(
    dimension: u16,
    vectors: &[f32],
    cell_ids: &[Id],
    graph_degree: usize,
) -> Result<CagraCudaBuildShape, CagraError>

pub fn validate_cuda_nndescent_build_inputs(
    dimension: u16,
    vectors: &[f32],
    cell_ids: &[Id],
    intermediate_graph_degree: usize,
    graph_degree: usize,
) -> Result<CagraCudaBuildShape, CagraError>

pub fn estimate_cuda_build_device_bytes(
    shape: CagraCudaBuildShape,
) -> Result<usize, CagraError>

pub fn validate_cuda_build_device_memory_budget(
    shape: CagraCudaBuildShape,
    device_memory_budget_bytes: Option<u64>,
    context: &str,
) -> Result<(), CagraError>

pub fn checked_cuda_build_flat_len(
    row_count: usize,
    row_width: usize,
    context: &str,
) -> Result<usize, CagraError>

pub fn decode_cuda_build_adjacency_rows(
    flat_adjacency: Vec<u32>,
    graph_degree: usize,
) -> Result<Vec<Vec<u32>>, CagraError>

Implementation rule: paste the current function bodies unchanged except for imports and visibility. This keeps exact-build and NN-descent validation behavior stable.

  • Step 3: Move existing resident-cache helpers into cuda_common

Move these exact functions from src/apps/cagra/gpu/cuda/resident.rs into cuda_common.rs and keep their signatures:

pub fn estimate_segment_device_bytes(
    row_count: u32,
    dimension: u32,
    graph_degree: u32,
) -> Option<usize>

pub fn select_resident_cache_evictions_to_fit<K>(
    resident_bytes: &HashMap<K, usize>,
    lru_oldest_to_newest: &[K],
    incoming_key: K,
    incoming_bytes: usize,
    capacity_bytes: Option<u64>,
    context: &str,
) -> Result<Vec<K>, CagraError>
where
    K: Copy + Eq + Hash

Implementation rule: paste the current function bodies unchanged except for imports and visibility.

  • Step 4: Add shared search launch helper

Move cagra_cuda_search_launch_resources, highest_power_of_two_at_most, bounded_output_count, checked_bytes, and usize_to_u32 from src/apps/cagra/gpu/cuda/search.rs into cuda_common.rs.

Use these public signatures:

pub fn cagra_cuda_search_launch_resources(
    row_count: usize,
    graph_degree: usize,
    config: CagraSearchConfig,
    dynamic_shared_budget_bytes: usize,
) -> Result<CagraCudaSearchLaunchResources, CagraError>

pub fn bounded_output_count(k: usize, row_count: usize) -> usize

pub fn checked_bytes(
    element_count: usize,
    element_size: usize,
    label: &str,
) -> Result<usize, CagraError>

pub fn usize_to_u32(value: usize, field_name: &str) -> Result<u32, CagraError>
  • Step 5: Add descriptor key helper

Add this helper to cuda_common.rs:

pub fn cagra_cuda_segment_key_from_descriptor(
    schema_id: u32,
    field_id: u64,
    descriptor: &CagraSegmentDescriptor,
    preferred_device_ordinal: Option<u32>,
) -> CagraCudaSegmentKey {
    CagraCudaSegmentKey {
        device_ordinal: preferred_device_ordinal.unwrap_or(0),
        schema_id,
        field_id,
        segment_id: descriptor.segment_id,
        generation: descriptor.generation,
        checksum: descriptor.checksum,
    }
}
  • Step 6: Export cuda_common

Modify src/apps/cagra/gpu/mod.rs:

pub mod config;
pub mod cuda_common;
pub mod trace;

#[cfg(feature = "cagra-cuda")]
pub mod cuda_oxide;

#[cfg(feature = "cagra-cuda-legacy")]
#[path = "cuda/mod.rs"]
pub mod cuda_legacy;

#[cfg(feature = "cagra-cuda")]
pub mod cuda {
    pub use super::cuda_oxide::{build, delta, resident, runtime, search};
}

#[cfg(not(feature = "cagra-cuda"))]
pub mod cuda {
    pub mod build {
        pub use super::super::cuda_common::{
            checked_cuda_build_flat_len, decode_cuda_build_adjacency_rows,
            estimate_cuda_build_device_bytes, select_cuda_build_algorithm,
            validate_cuda_build_device_memory_budget, validate_cuda_build_inputs,
            validate_cuda_nndescent_build_inputs, CagraCudaBuildShape, CudaCagraBuildAlgorithm,
            CAGRA_CUDA_BUILD_AUTO_EXACT_ROW_THRESHOLD, CAGRA_CUDA_BUILD_MAX_GRAPH_DEGREE,
        };
    }
}

#[cfg(feature = "cagra-vulkan")]
pub mod vulkan;
  • Step 7: Update imports

Change old and oxide modules to import from crate::apps::cagra::gpu::cuda_common.

Required import shape in src/apps/cagra/gpu/cuda_oxide/build.rs:

use crate::apps::cagra::{
    gpu::cuda_common::{
        checked_cuda_build_flat_len, decode_cuda_build_adjacency_rows,
        validate_cuda_build_device_memory_budget, validate_cuda_build_inputs, CagraCudaBuildShape,
        CAGRA_CUDA_BUILD_MAX_GRAPH_DEGREE,
    },
    search::validated_base_segment_layout,
    types::{CagraBaseSegment, CagraError, CAGRA_ARTIFACT_VERSION},
};

Required import shape in temporary legacy modules:

use crate::apps::cagra::gpu::cuda_common::{
    checked_cuda_build_flat_len, decode_cuda_build_adjacency_rows,
    validate_cuda_build_device_memory_budget, validate_cuda_build_inputs,
    validate_cuda_nndescent_build_inputs, CagraCudaBuildShape, CudaCagraBuildAlgorithm,
    CAGRA_CUDA_BUILD_AUTO_EXACT_ROW_THRESHOLD, CAGRA_CUDA_BUILD_MAX_GRAPH_DEGREE,
};
  • Step 8: Run host-only tests

Run:

cargo test --lib cuda_build_device_budget
cargo test --lib cuda_nndescent_decode
cargo test --lib cuda_nndescent_build_length_helper

Expected: all moved helper tests pass.

  • Step 9: Commit

Run:

git add src/apps/cagra/gpu/cuda_common.rs src/apps/cagra/gpu/mod.rs src/apps/cagra/gpu/cuda src/apps/cagra/gpu/cuda_oxide
git commit -m "refactor(cagra): split CUDA host helpers"

Task 3: Cut Feature Flags To cuda-oxide With Legacy Oracle

Files:

  • Modify: Cargo.toml

  • Modify: src/bin/cagra_cuda_oxide_smoke.rs

  • Modify: src/bin/cagra_cuda_oxide_build_smoke.rs

  • Step 1: Update features

Replace the CAGRA CUDA feature section in Cargo.toml with:

# Enable CAGRA custom CUDA kernels through cuda-oxide.
# Build and run cuda-oxide binaries with `cargo oxide`.
cagra-cuda = ["dep:cuda-device", "dep:cuda-host", "dep:cuda-core"]
# Temporary parity oracle for the previous cudarc/NVRTC CAGRA CUDA path.
# This feature must not be used by production service routing.
cagra-cuda-legacy = ["dep:cudarc"]
# Backward-compatible alias for existing smoke commands during migration.
cagra-cuda-oxide = ["cagra-cuda"]
  • Step 2: Update smoke binary required features

Gate both existing smoke binaries by the new production CUDA feature. Keep cagra-cuda-oxide only as an alias for old invocations during migration; it must not be the required feature for new smoke targets because the required verification commands use --features cagra-cuda.

[[bin]]
name = "cagra_cuda_oxide_smoke"
path = "src/bin/cagra_cuda_oxide_smoke.rs"
required-features = ["cagra-cuda"]

[[bin]]
name = "cagra_cuda_oxide_build_smoke"
path = "src/bin/cagra_cuda_oxide_build_smoke.rs"
required-features = ["cagra-cuda"]
  • Step 3: Change legacy comparisons to use cuda_legacy

In src/bin/cagra_cuda_oxide_smoke.rs, replace:

#[cfg(feature = "cagra-cuda")]
{
    let cuda_hits = morpheus::apps::cagra::gpu::cuda::delta::search_delta_segment_cuda(

with:

#[cfg(feature = "cagra-cuda-legacy")]
{
    let cuda_hits = morpheus::apps::cagra::gpu::cuda_legacy::delta::search_delta_segment_cuda(

In src/bin/cagra_cuda_oxide_build_smoke.rs, replace:

#[cfg(feature = "cagra-cuda")]
{
    let cuda = morpheus::apps::cagra::gpu::cuda::build::build_base_segment_exact_l2_cuda(

with:

#[cfg(feature = "cagra-cuda-legacy")]
{
    let cuda = morpheus::apps::cagra::gpu::cuda_legacy::build::build_base_segment_exact_l2_cuda(
  • Step 4: Run feature smoke without legacy

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_build_smoke --features cagra-cuda

Expected output contains:

PASSED: cuda-oxide CAGRA exact build smoke matched CPU exact-L2 oracle
  • Step 5: Run feature smoke with legacy oracle

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_build_smoke --features cagra-cuda,cagra-cuda-legacy

Expected output contains:

PASSED: cuda-oxide CAGRA exact build smoke matched existing .cu path
  • Step 6: Commit

Run:

git add Cargo.toml src/bin/cagra_cuda_oxide_smoke.rs src/bin/cagra_cuda_oxide_build_smoke.rs
git commit -m "build(cagra): make cuda-oxide the CUDA feature"

Task 4: Add cuda-oxide Runtime And Resident Cache

Files:

  • Create: src/apps/cagra/gpu/cuda_oxide/runtime.rs

  • Create: src/apps/cagra/gpu/cuda_oxide/resident.rs

  • Modify: src/apps/cagra/gpu/cuda_oxide/mod.rs

  • Test: inline tests in resident.rs.

  • Step 1: Export modules

Modify src/apps/cagra/gpu/cuda_oxide/mod.rs:

pub mod build;
pub mod delta;
pub mod resident;
pub mod runtime;
pub mod search;
  • Step 2: Add runtime API

Create src/apps/cagra/gpu/cuda_oxide/runtime.rs:

use std::{
    collections::HashMap,
    sync::{Arc, Mutex, OnceLock},
};

use cuda_core::CudaContext;

use crate::apps::cagra::CagraError;

const CONTEXT_CACHE_POISON_MESSAGE: &str = "CAGRA cuda-oxide context cache mutex poisoned";

static CAGRA_CUDA_OXIDE_CONTEXT_CACHE: OnceLock<Mutex<HashMap<u32, Arc<CudaContext>>>> =
    OnceLock::new();

#[derive(Debug, Clone, PartialEq, Eq)]
pub struct CagraCudaRuntimeProbe {
    pub feature_compiled: bool,
    pub available: bool,
    pub device_ordinal: Option<u32>,
    pub reason: String,
}

pub fn probe_cagra_cuda_runtime(preferred_device_ordinal: Option<u32>) -> CagraCudaRuntimeProbe {
    let device_ordinal = preferred_device_ordinal.unwrap_or(0);
    match CudaContext::new(device_ordinal as usize) {
        Ok(_) => CagraCudaRuntimeProbe {
            feature_compiled: true,
            available: true,
            device_ordinal: Some(device_ordinal),
            reason: format!(
                "cagra-cuda feature compiled with cuda-oxide; CUDA context initialized for device ordinal {device_ordinal}"
            ),
        },
        Err(error) => CagraCudaRuntimeProbe {
            feature_compiled: true,
            available: false,
            device_ordinal: Some(device_ordinal),
            reason: format!(
                "cagra-cuda feature compiled with cuda-oxide; CUDA context unavailable for device ordinal {device_ordinal}: {error}"
            ),
        },
    }
}

pub fn cagra_cuda_context(
    preferred_device_ordinal: Option<u32>,
) -> Result<Arc<CudaContext>, CagraError> {
    let device_ordinal = preferred_device_ordinal.unwrap_or(0);
    let cache = CAGRA_CUDA_OXIDE_CONTEXT_CACHE.get_or_init(Default::default);
    if let Some(context) = cache
        .lock()
        .expect(CONTEXT_CACHE_POISON_MESSAGE)
        .get(&device_ordinal)
        .cloned()
    {
        return Ok(context);
    }

    let context = CudaContext::new(device_ordinal as usize).map_err(|error| {
        CagraError::Storage(format!(
            "failed to initialize CAGRA cuda-oxide context for device ordinal {device_ordinal}: {error}"
        ))
    })?;
    let mut contexts = cache.lock().expect(CONTEXT_CACHE_POISON_MESSAGE);
    if let Some(existing) = contexts.get(&device_ordinal).cloned() {
        return Ok(existing);
    }
    contexts.insert(device_ordinal, Arc::clone(&context));
    Ok(context)
}

pub(crate) fn clear_cagra_cuda_runtime_caches() {
    let cleared_contexts = {
        let cache = CAGRA_CUDA_OXIDE_CONTEXT_CACHE.get_or_init(Default::default);
        let mut contexts = cache.lock().expect(CONTEXT_CACHE_POISON_MESSAGE);
        std::mem::take(&mut *contexts)
    };
    drop(cleared_contexts);
}
  • Step 3: Add resident segment API

Create src/apps/cagra/gpu/cuda_oxide/resident.rs by porting the current src/apps/cagra/gpu/cuda/resident.rs structure with these concrete type changes:

use std::{
    collections::{HashMap, VecDeque},
    sync::{Arc, Mutex, MutexGuard},
};

use cuda_core::{CudaContext, DeviceBuffer};

use crate::apps::cagra::{
    gpu::cuda_common::{estimate_segment_device_bytes, select_resident_cache_evictions_to_fit},
    layout::CAGRA_EMPTY_NEIGHBOR,
    CagraError,
};

const CACHE_POISON_MESSAGE: &str = "CAGRA cuda-oxide resident segment cache mutex poisoned";

pub use crate::apps::cagra::gpu::cuda_common::CagraCudaSegmentKey;

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub struct CagraCudaResidentSegmentMeta {
    pub key: CagraCudaSegmentKey,
    pub row_count: u32,
    pub dimension: u32,
    pub graph_degree: u32,
    pub device_bytes: usize,
}

pub struct CagraCudaResidentSegment {
    meta: CagraCudaResidentSegmentMeta,
    context: Arc<CudaContext>,
    vectors: DeviceBuffer<f32>,
    adjacency: DeviceBuffer<u32>,
}

Port these methods from legacy without changing external behavior:

impl CagraCudaResidentSegment {
    pub fn meta(&self) -> CagraCudaResidentSegmentMeta
    pub fn key(&self) -> CagraCudaSegmentKey
    pub fn row_count(&self) -> u32
    pub fn dimension(&self) -> u32
    pub fn graph_degree(&self) -> u32
    pub fn device_bytes(&self) -> usize
    pub fn context(&self) -> &Arc<CudaContext>
    pub fn vectors(&self) -> &DeviceBuffer<f32>
    pub fn adjacency(&self) -> &DeviceBuffer<u32>
}
  • Step 4: Port upload with cuda-oxide buffer APIs

In upload_resident_segment, replace cudarc copies:

let stream = context.default_stream();
let vectors = DeviceBuffer::from_host(&stream, vectors).map_err(|error| {
    CagraError::Storage(format!(
        "failed to upload CAGRA cuda-oxide resident segment vectors for segment {}: {error}",
        key.segment_id
    ))
})?;
let adjacency = DeviceBuffer::from_host(&stream, flat_adjacency).map_err(|error| {
    CagraError::Storage(format!(
        "failed to upload CAGRA cuda-oxide resident segment adjacency for segment {}: {error}",
        key.segment_id
    ))
})?;
stream.synchronize().map_err(|error| {
    CagraError::Storage(format!(
        "failed to synchronize CAGRA cuda-oxide resident upload for segment {}: {error}",
        key.segment_id
    ))
})?;
  • Step 5: Port cache type unchanged

Port CagraCudaResidentSegmentCache, CagraCudaResidentSegmentCacheState, evict_state_to_fit, evict_state_to_target_bytes, mark_recent, and remove_lru_entry from legacy. Keep public method names unchanged:

pub fn get(&self, key: CagraCudaSegmentKey) -> Option<Arc<CagraCudaResidentSegment>>
pub fn insert(&self, segment: CagraCudaResidentSegment) -> Arc<CagraCudaResidentSegment>
pub fn insert_with_capacity(
    &self,
    segment: CagraCudaResidentSegment,
    capacity_bytes: Option<u64>,
) -> Result<Arc<CagraCudaResidentSegment>, CagraError>
pub fn evict_to_fit(
    &self,
    incoming_key: CagraCudaSegmentKey,
    incoming_bytes: usize,
    capacity_bytes: Option<u64>,
    context: &str,
) -> Result<(), CagraError>
pub fn evict_to_target_bytes(&self, target_bytes: usize)
pub fn remove(&self, key: CagraCudaSegmentKey) -> Option<Arc<CagraCudaResidentSegment>>
pub fn clear(&self)
pub fn len(&self) -> usize
pub fn resident_bytes(&self) -> usize
  • Step 6: Run host-only resident tests

Run:

cargo test --lib select_resident_cache_evictions_to_fit

Expected: all resident cache eviction tests pass.

  • Step 7: Run oxide compile smoke

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_build_smoke --features cagra-cuda

Expected output contains:

PASSED: cuda-oxide CAGRA exact build smoke matched CPU exact-L2 oracle
  • Step 8: Commit

Run:

git add src/apps/cagra/gpu/cuda_oxide/runtime.rs src/apps/cagra/gpu/cuda_oxide/resident.rs src/apps/cagra/gpu/cuda_oxide/mod.rs
git commit -m "feat(cagra): add cuda-oxide runtime and resident cache"

Task 5: Port Base Graph Search To cuda-oxide

Files:

  • Create: src/apps/cagra/gpu/cuda_oxide/search.rs

  • Create: src/bin/cagra_cuda_oxide_search_smoke.rs

  • Modify: Cargo.toml

  • Test: src/bin/cagra_cuda_oxide_search_smoke.rs

  • Step 1: Add search module host shell

Create src/apps/cagra/gpu/cuda_oxide/search.rs with these imports and public functions:

use std::{
    cmp::Ordering,
    collections::BTreeMap,
    mem::size_of,
    sync::{Arc, OnceLock},
};

use cuda_core::{DeviceBuffer, LaunchConfig};
use cuda_host::cuda_module;
use dovahkiin::types::Id;

pub use crate::apps::cagra::backend::{
    CAGRA_CUDA_SEARCH_MAX_GRAPH_DEGREE, CAGRA_CUDA_SEARCH_MAX_ITERATIONS,
    CAGRA_CUDA_SEARCH_MAX_K, CAGRA_CUDA_SEARCH_MAX_ROWS, CAGRA_CUDA_SEARCH_MAX_WIDTH,
};

use crate::apps::cagra::{
    artifact::{artifact_checksum, encode_artifact},
    backend::{
        estimate_cagra_cuda_base_search_device_bytes as estimate_base_search_device_bytes,
        estimate_cagra_cuda_resident_device_bytes, validate_cagra_cuda_cache_capacity_bytes,
        validate_cagra_cuda_device_memory_budget_bytes, CagraBackendRoute,
        CagraBackendSearchResult, CagraCudaSearchShape, CagraExecutionBackendKind,
    },
    gpu::{
        cuda_common::{
            bounded_output_count, cagra_cuda_search_launch_resources,
            cagra_cuda_segment_key_from_descriptor, checked_bytes, usize_to_u32,
            CagraCudaSegmentKey, CAGRA_CUDA_DEFAULT_DYNAMIC_SHARED_BUDGET_BYTES,
            CAGRA_CUDA_EMPTY_OUTPUT_ROW, CAGRA_CUDA_SEARCH_BLOCK_THREADS,
        },
        trace::CagraGpuSearchTrace,
    },
    layout::FixedDegreeGraphLayout,
    search::{validate_base_segment_query, validated_base_segment_layout},
    types::{
        CagraArtifactKind, CagraBaseSegment, CagraError, CagraSearchConfig, CagraSearchHit,
        CagraSegmentDescriptor, CagraSegmentKind,
    },
};

use super::{
    resident::{
        upload_resident_segment, CagraCudaResidentSegment, CagraCudaResidentSegmentCache,
    },
    runtime::{cagra_cuda_context, clear_cagra_cuda_runtime_caches},
};

static RESIDENT_SEGMENT_CACHE: OnceLock<CagraCudaResidentSegmentCache> = OnceLock::new();

#[derive(Debug, Clone, Copy)]
struct CagraCudaSearchRowResult {
    row_ordinal: u32,
    distance: f32,
}
  • Step 2: Port host validation and cache preparation

Copy these functions from legacy src/apps/cagra/gpu/cuda/search.rs into the oxide file and update error text from CAGRA CUDA to CAGRA cuda-oxide only when the message names implementation:

pub fn search_base_segment_graph_cuda(
    schema_id: u32,
    field_id: u64,
    descriptor: &CagraSegmentDescriptor,
    segment: &CagraBaseSegment,
    query: &[f32],
    config: CagraSearchConfig,
    preferred_device_ordinal: Option<u32>,
) -> Result<CagraBackendSearchResult, CagraError>

pub(crate) fn search_base_segment_graph_cuda_with_budgets(
    schema_id: u32,
    field_id: u64,
    descriptor: &CagraSegmentDescriptor,
    segment: &CagraBaseSegment,
    query: &[f32],
    config: CagraSearchConfig,
    preferred_device_ordinal: Option<u32>,
    device_memory_budget_bytes: Option<u64>,
    cache_capacity_bytes: Option<u64>,
) -> Result<CagraBackendSearchResult, CagraError>

Use cuda_common::cagra_cuda_segment_key_from_descriptor instead of a local duplicate. Use cuda_common::checked_bytes and cuda_common::bounded_output_count.

  • Step 3: Add kernel translation header

Add this #[cuda_module] shell above the host functions:

const CAGRA_ROW_STATE_UNSEEN: u32 = 0;
const CAGRA_ROW_STATE_QUEUED: u32 = 1;
const CAGRA_ROW_STATE_EXPANDED: u32 = 2;
const FLOAT_MAX: f32 = 3.402823466e38_f32;

#[cuda_module]
mod kernels {
    use cuda_device::{kernel, thread, DisjointSlice, DynamicSharedArray};

    use super::{
        CAGRA_ROW_STATE_EXPANDED, CAGRA_ROW_STATE_QUEUED, CAGRA_ROW_STATE_UNSEEN,
        CAGRA_CUDA_EMPTY_OUTPUT_ROW, CAGRA_CUDA_SEARCH_BLOCK_THREADS, FLOAT_MAX,
    };

    #[kernel]
    pub fn cagra_search_l2_oxide(
        vectors: &[f32],
        adjacency: &[u32],
        query: &[f32],
        row_count: u32,
        dimension: u32,
        graph_degree: u32,
        k: u32,
        search_width: u32,
        max_iterations: u32,
        hash_table_slots: u32,
        output_capacity: u32,
        mut frontier_rows: DisjointSlice<u32>,
        mut frontier_distances: DisjointSlice<f32>,
        mut row_states: DisjointSlice<u32>,
        mut output_rows: DisjointSlice<u32>,
        mut output_distances: DisjointSlice<f32>,
    ) {
        let hash_table: *mut u32 = DynamicSharedArray::<u32>::get();
        let tid = thread::threadIdx_x();
        if thread::blockIdx_x() != 0
            || thread::blockDim_x() != CAGRA_CUDA_SEARCH_BLOCK_THREADS
            || hash_table_slots == 0
        {
            return;
        }

        let final_capacity = min_u32(min_u32(output_capacity, k), row_count);
        initialize_search_state(
            tid,
            final_capacity,
            row_count,
            hash_table_slots,
            hash_table,
            &mut output_rows,
            &mut output_distances,
            &mut row_states,
        );
        thread::sync_threads();

        if row_count == 0 || final_capacity == 0 {
            return;
        }

        run_single_cta_graph_search(
            vectors,
            adjacency,
            query,
            row_count,
            dimension,
            graph_degree,
            search_width,
            max_iterations,
            hash_table_slots,
            final_capacity,
            hash_table,
            &mut frontier_rows,
            &mut frontier_distances,
            &mut row_states,
            &mut output_rows,
            &mut output_distances,
        );
    }
}
  • Step 4: Port device helpers mechanically

Inside mod kernels, translate each helper from src/apps/cagra/gpu/cuda/kernels/search_l2.cu into Rust with this mapping:

cagra_distance_row_less -> fn distance_row_less with four scalar arguments and bool return
cagra_squared_l2_row_cooperative -> fn squared_l2_row_cooperative with vectors, query, row ordinal, dimension, and partial sums
cagra_frontier_insert -> fn frontier_insert with frontier slices, length pointer, capacity, row ordinal, and distance
cagra_frontier_pop_front -> fn frontier_pop_front with frontier slices and length pointer
cagra_output_insert -> fn output_insert with output slices, capacity, output length pointer, row ordinal, and distance
cagra_hash_contains -> fn hash_contains with hash table pointer, slot count, and row ordinal
cagra_hash_insert -> fn hash_insert with hash table pointer, slot count, and row ordinal
cagra_hash_reset_and_reseed -> fn hash_reset_and_reseed with hash table, frontier slices, output slices, and active lengths

Use these exact cuda-oxide translation rules:

threadIdx.x -> thread::threadIdx_x()
blockIdx.x -> thread::blockIdx_x()
blockDim.x -> thread::blockDim_x()
__syncthreads() -> thread::sync_threads()
extern __shared__ unsigned int hash_table[] -> DynamicSharedArray::<u32>::get()
unsigned int* output_rows -> &mut DisjointSlice<u32>
float* output_distances -> &mut DisjointSlice<f32>
array[index] = value -> unsafe { *slice.get_unchecked_mut(index as usize) = value; }
array[index] read from DisjointSlice -> unsafe { *slice.get_unchecked_mut(index as usize) }
hash_table[slot] -> unsafe { *hash_table.add(slot as usize) }

The port must preserve tie-breaking: lower distance wins, then lower row ordinal wins.

  • Step 5: Implement oxide launch wrapper

In cagra_cuda_search_launch, replace cudarc launch builder with cuda-oxide typed launch:

let stream = resident.context().default_stream();
let d_query = DeviceBuffer::from_host(&stream, query).map_err(|error| {
    CagraError::Storage(format!(
        "failed to upload CAGRA cuda-oxide query for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let launch_resources = cagra_cuda_search_launch_resources(
    row_count,
    usize::try_from(graph_degree_u32).expect("graph degree should fit usize"),
    config,
    CAGRA_CUDA_DEFAULT_DYNAMIC_SHARED_BUDGET_BYTES,
)?;
let output_count = launch_resources.output_count;
let mut d_output_rows = DeviceBuffer::<u32>::zeroed(&stream, output_count).map_err(|error| {
    CagraError::Storage(format!(
        "failed to allocate CAGRA cuda-oxide output rows for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let mut d_output_distances = DeviceBuffer::<f32>::zeroed(&stream, output_count).map_err(|error| {
    CagraError::Storage(format!(
        "failed to allocate CAGRA cuda-oxide output distances for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let mut d_frontier_rows = DeviceBuffer::<u32>::zeroed(&stream, row_count).map_err(|error| {
    CagraError::Storage(format!(
        "failed to allocate CAGRA cuda-oxide frontier rows for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let mut d_frontier_distances = DeviceBuffer::<f32>::zeroed(&stream, row_count).map_err(|error| {
    CagraError::Storage(format!(
        "failed to allocate CAGRA cuda-oxide frontier distances for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let mut d_row_states = DeviceBuffer::<u32>::zeroed(&stream, row_count).map_err(|error| {
    CagraError::Storage(format!(
        "failed to allocate CAGRA cuda-oxide row-state buffer for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let module = kernels::load(resident.context()).map_err(|error| {
    CagraError::Storage(format!(
        "failed to load embedded CAGRA cuda-oxide search module for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
module
    .cagra_search_l2_oxide(
        &stream,
        LaunchConfig {
            grid_dim: (1, 1, 1),
            block_dim: (CAGRA_CUDA_SEARCH_BLOCK_THREADS, 1, 1),
            shared_mem_bytes: launch_resources.shared_mem_bytes,
        },
        resident.vectors(),
        resident.adjacency(),
        &d_query,
        row_count_u32,
        dimension_u32,
        graph_degree_u32,
        k_u32,
        search_width_u32,
        max_iterations_u32,
        hash_table_slots_u32,
        output_capacity_u32,
        &mut d_frontier_rows,
        &mut d_frontier_distances,
        &mut d_row_states,
        &mut d_output_rows,
        &mut d_output_distances,
    )
    .map_err(|error| {
        CagraError::Storage(format!(
            "failed to launch CAGRA cuda-oxide search kernel for segment {}: {error}",
            resident.key().segment_id
        ))
    })?;
stream.synchronize().map_err(|error| {
    CagraError::Storage(format!(
        "failed to synchronize CAGRA cuda-oxide search kernel for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let host_rows = d_output_rows.to_host_vec(&stream).map_err(|error| {
    CagraError::Storage(format!(
        "failed to copy CAGRA cuda-oxide output rows for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
let host_distances = d_output_distances.to_host_vec(&stream).map_err(|error| {
    CagraError::Storage(format!(
        "failed to copy CAGRA cuda-oxide output distances for segment {}: {error}",
        resident.key().segment_id
    ))
})?;
  • Step 6: Add search route reason

Use this route function:

fn cagra_cuda_route() -> CagraBackendRoute {
    CagraBackendRoute {
        selected: CagraExecutionBackendKind::CudaCustom,
        reason: "admitted to CAGRA cuda-oxide device-resident graph search kernel v1".to_string(),
    }
}
  • Step 7: Add search smoke binary

Add to Cargo.toml:

[[bin]]
name = "cagra_cuda_oxide_search_smoke"
path = "src/bin/cagra_cuda_oxide_search_smoke.rs"
required-features = ["cagra-cuda"]

Create src/bin/cagra_cuda_oxide_search_smoke.rs:

use morpheus::apps::cagra::{
    artifact::{artifact_checksum, encode_artifact},
    gpu::cuda_oxide::{
        build::build_base_segment_exact_l2_cuda_oxide,
        search::search_base_segment_graph_cuda,
    },
    search::search_base_segment_graph_cpu,
    types::{
        CagraArtifactKind, CagraSearchConfig, CagraSegmentDescriptor, CagraSegmentKind,
        CAGRA_ARTIFACT_VERSION,
    },
};

fn main() {
    let dimension = 4u16;
    let vectors = vec![
        0.0, 0.0, 0.0, 0.0,
        1.0, 0.0, 0.0, 0.0,
        0.0, 1.0, 0.0, 0.0,
        0.0, 0.0, 1.0, 0.0,
        2.0, 0.0, 0.0, 0.0,
        0.0, 2.0, 0.0, 0.0,
    ];
    let cell_ids = (0..6).map(|value| dovahkiin::types::Id::new(value, 0)).collect::<Vec<_>>();
    let segment = build_base_segment_exact_l2_cuda_oxide(
        137,
        139,
        dimension,
        vectors.as_slice(),
        cell_ids.as_slice(),
        3,
        None,
        None,
    )
    .expect("cuda-oxide exact build should succeed");
    let encoded = encode_artifact(CagraArtifactKind::BaseSegment, &segment)
        .expect("base segment should encode");
    let descriptor = CagraSegmentDescriptor {
        artifact_version: CAGRA_ARTIFACT_VERSION,
        kind: CagraSegmentKind::Base,
        segment_id: segment.segment_id,
        generation: segment.generation,
        row_count: segment.cell_ids.len() as u64,
        checksum: artifact_checksum(encoded.as_slice()),
    };
    let query = vec![0.9, 0.0, 0.0, 0.0];
    let config = CagraSearchConfig {
        k: 3,
        search_width: 4,
        max_iterations: 16,
    };
    let cpu_hits = search_base_segment_graph_cpu(&segment, query.as_slice(), config)
        .expect("CPU search should succeed");
    let oxide_hits = search_base_segment_graph_cuda(
        17,
        19,
        &descriptor,
        &segment,
        query.as_slice(),
        config,
        None,
    )
    .expect("cuda-oxide graph search should succeed")
    .hits;
    if oxide_hits != cpu_hits {
        eprintln!("FAILED: cuda-oxide hits {oxide_hits:?} did not match CPU hits {cpu_hits:?}");
        std::process::exit(1);
    }
    println!("PASSED: cuda-oxide CAGRA base graph search smoke matched CPU oracle");
}
  • Step 8: Run search smoke

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_search_smoke --features cagra-cuda

Expected output contains:

PASSED: cuda-oxide CAGRA base graph search smoke matched CPU oracle
  • Step 9: Commit

Run:

git add Cargo.toml src/apps/cagra/gpu/cuda_oxide/search.rs src/bin/cagra_cuda_oxide_search_smoke.rs
git commit -m "feat(cagra): port base graph search to cuda-oxide"

Task 6: Port NN-descent Build To cuda-oxide

Files:

  • Modify: src/apps/cagra/gpu/cuda_oxide/build.rs

  • Create: src/bin/cagra_cuda_oxide_nndescent_smoke.rs

  • Modify: Cargo.toml

  • Step 1: Add NN-descent kernel constants

In src/apps/cagra/gpu/cuda_oxide/build.rs, add:

const CAGRA_CUDA_BUILD_NNDESCENT_REFINEMENT_ROUNDS: u32 = 4;
  • Step 2: Add public NN-descent build API

Add this public function to build.rs:

pub fn build_base_segment_nndescent_l2_cuda(
    segment_id: u64,
    generation: u64,
    dimension: u16,
    vectors: &[f32],
    cell_ids: &[Id],
    intermediate_graph_degree: usize,
    graph_degree: usize,
    preferred_device_ordinal: Option<u32>,
    device_memory_budget_bytes: Option<u64>,
) -> Result<CagraBaseSegment, CagraError> {
    let shape = validate_cuda_nndescent_build_inputs(
        dimension,
        vectors,
        cell_ids,
        intermediate_graph_degree,
        graph_degree,
    )?;
    validate_cuda_build_device_memory_budget(
        shape,
        device_memory_budget_bytes,
        "CAGRA cuda-oxide NN-descent build",
    )?;
    let intermediate_adjacency = if shape.row_count == 0 {
        Vec::new()
    } else {
        decode_cuda_build_adjacency_rows(
            launch_cagra_build_nndescent_l2_oxide(
                vectors,
                shape,
                segment_id,
                preferred_device_ordinal,
            )?,
            intermediate_graph_degree,
        )?
    };
    let adjacency = optimize_cagra_intermediate_graph_cpu(
        vectors,
        usize::from(dimension),
        intermediate_adjacency.as_slice(),
        intermediate_graph_degree,
        graph_degree,
    )?;
    let segment = CagraBaseSegment {
        artifact_version: CAGRA_ARTIFACT_VERSION,
        segment_id,
        generation,
        dimension,
        vectors: vectors.to_vec(),
        cell_ids: cell_ids.to_vec(),
        adjacency,
    };
    let _ = validated_base_segment_layout(&segment)?;
    Ok(segment)
}
  • Step 3: Port kernel

In the existing #[cuda_module] mod kernels, translate src/apps/cagra/gpu/cuda/kernels/build_nndescent_l2.cu into:

#[kernel]
pub fn cagra_build_nndescent_l2_oxide(
    vectors: &[f32],
    row_count: u32,
    dimension: u32,
    graph_degree: u32,
    refinement_rounds: u32,
    mut out_adjacency: DisjointSlice<u32>,
) {
    let row = thread::blockIdx_x();
    if row >= row_count || thread::threadIdx_x() != 0 {
        return;
    }
    if graph_degree == 0 || graph_degree as usize > MAX_GRAPH_DEGREE {
        return;
    }

    initialize_deterministic_neighbors(row, row_count, graph_degree, &mut out_adjacency);
    let mut round = 0;
    while round < refinement_rounds {
        refine_row_neighbors_l2(row, vectors, row_count, dimension, graph_degree, &mut out_adjacency);
        round += 1;
    }
    sort_row_neighbors_l2(row, vectors, row_count, dimension, graph_degree, &mut out_adjacency);
}

Mechanical port rule: keep the current .cu deterministic initialization, distance scoring, neighbor replacement, duplicate avoidance, and final row sorting behavior. The function names above are fixed so tests can inspect generated PTX names.

  • Step 4: Add NN-descent launch wrapper

Add:

fn launch_cagra_build_nndescent_l2_oxide(
    vectors: &[f32],
    shape: CagraCudaBuildShape,
    segment_id: u64,
    preferred_device_ordinal: Option<u32>,
) -> Result<Vec<u32>, CagraError> {
    let row_count = shape.row_count as usize;
    let graph_degree = shape.graph_degree as usize;
    let output_len = checked_cuda_build_flat_len(row_count, graph_degree, "adjacency")?;
    let device_ordinal = usize::try_from(preferred_device_ordinal.unwrap_or(0)).map_err(|_| {
        CagraError::InvalidConfig("cuda-oxide preferred device ordinal does not fit usize".into())
    })?;
    let ctx = CudaContext::new(device_ordinal).map_err(|error| {
        CagraError::Storage(format!(
            "failed to create cuda-oxide context for device {device_ordinal}: {error}"
        ))
    })?;
    let stream = ctx.default_stream();
    let d_vectors = DeviceBuffer::from_host(&stream, vectors).map_err(|error| {
        CagraError::Storage(format!(
            "failed to upload CAGRA cuda-oxide NN-descent vectors for segment {segment_id}: {error}"
        ))
    })?;
    let mut d_adjacency = DeviceBuffer::<u32>::zeroed(&stream, output_len).map_err(|error| {
        CagraError::Storage(format!(
            "failed to allocate CAGRA cuda-oxide NN-descent adjacency for segment {segment_id}: {error}"
        ))
    })?;
    let module = kernels::load(&ctx).map_err(|error| {
        CagraError::Storage(format!(
            "failed to load embedded cuda-oxide NN-descent module for segment {segment_id}: {error}"
        ))
    })?;
    module
        .cagra_build_nndescent_l2_oxide(
            &stream,
            LaunchConfig {
                grid_dim: (shape.row_count, 1, 1),
                block_dim: (1, 1, 1),
                shared_mem_bytes: 0,
            },
            &d_vectors,
            shape.row_count,
            shape.dimension,
            shape.graph_degree,
            CAGRA_CUDA_BUILD_NNDESCENT_REFINEMENT_ROUNDS,
            &mut d_adjacency,
        )
        .map_err(|error| {
            CagraError::Storage(format!(
                "failed to launch CAGRA cuda-oxide NN-descent build kernel for segment {segment_id}: {error}"
            ))
        })?;
    stream.synchronize().map_err(|error| {
        CagraError::Storage(format!(
            "failed to synchronize CAGRA cuda-oxide NN-descent build kernel for segment {segment_id}: {error}"
        ))
    })?;
    d_adjacency.to_host_vec(&stream).map_err(|error| {
        CagraError::Storage(format!(
            "failed to copy CAGRA cuda-oxide NN-descent adjacency for segment {segment_id}: {error}"
        ))
    })
}
  • Step 5: Add NN-descent smoke

Add to Cargo.toml:

[[bin]]
name = "cagra_cuda_oxide_nndescent_smoke"
path = "src/bin/cagra_cuda_oxide_nndescent_smoke.rs"
required-features = ["cagra-cuda"]

Create src/bin/cagra_cuda_oxide_nndescent_smoke.rs:

use morpheus::apps::cagra::{
    gpu::cuda_oxide::build::build_base_segment_nndescent_l2_cuda,
    search::validated_base_segment_layout,
};

fn main() {
    let dimension = 4u16;
    let mut vectors = Vec::new();
    for row in 0..32 {
        vectors.extend_from_slice(&[
            row as f32,
            (row % 7) as f32,
            (row % 5) as f32,
            (row % 3) as f32,
        ]);
    }
    let cell_ids = (0..32).map(|value| dovahkiin::types::Id::new(value, 0)).collect::<Vec<_>>();
    let segment = build_base_segment_nndescent_l2_cuda(
        211,
        223,
        dimension,
        vectors.as_slice(),
        cell_ids.as_slice(),
        8,
        4,
        None,
        None,
    )
    .expect("cuda-oxide NN-descent build should succeed");
    let layout = validated_base_segment_layout(&segment).expect("NN-descent layout should validate");
    assert_eq!(layout.row_count(), 32);
    assert_eq!(layout.graph_degree(), 4);
    println!("PASSED: cuda-oxide CAGRA NN-descent build smoke produced a valid fixed-degree graph");
}
  • Step 6: Run NN-descent smoke

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_nndescent_smoke --features cagra-cuda

Expected output contains:

PASSED: cuda-oxide CAGRA NN-descent build smoke produced a valid fixed-degree graph
  • Step 7: Commit

Run:

git add Cargo.toml src/apps/cagra/gpu/cuda_oxide/build.rs src/bin/cagra_cuda_oxide_nndescent_smoke.rs
git commit -m "feat(cagra): port NN-descent build to cuda-oxide"

Task 7: Cut CAGRA Service Routing Over To cuda-oxide

Files:

  • Modify: src/apps/cagra/service.rs

  • Modify: src/apps/cagra/tests.rs

  • Create: src/bin/cagra_cuda_oxide_e2e_smoke.rs

  • Modify: Cargo.toml

  • Step 1: Keep service imports stable

Because gpu::cuda now re-exports cuda_oxide, most service call sites should keep the same paths:

super::gpu::cuda::runtime::probe_cagra_cuda_runtime
super::gpu::cuda::search::search_base_segment_graph_cuda_with_budgets
super::gpu::cuda::delta::search_delta_segment_cuda_with_budgets
super::gpu::cuda::build::build_base_segment_exact_l2_cuda
super::gpu::cuda::build::build_base_segment_nndescent_l2_cuda

Required source behavior: those paths must resolve to cuda-oxide modules with --features cagra-cuda.

  • Step 2: Add exact-build compatibility function

In src/apps/cagra/gpu/cuda_oxide/build.rs, expose the production name as a wrapper:

pub fn build_base_segment_exact_l2_cuda(
    segment_id: u64,
    generation: u64,
    dimension: u16,
    vectors: &[f32],
    cell_ids: &[Id],
    graph_degree: usize,
    preferred_device_ordinal: Option<u32>,
    device_memory_budget_bytes: Option<u64>,
) -> Result<CagraBaseSegment, CagraError> {
    build_base_segment_exact_l2_cuda_oxide(
        segment_id,
        generation,
        dimension,
        vectors,
        cell_ids,
        graph_degree,
        preferred_device_ordinal,
        device_memory_budget_bytes,
    )
}
  • Step 3: Add delta compatibility function

In src/apps/cagra/gpu/cuda_oxide/delta.rs, expose the production name as a wrapper:

pub fn search_delta_segment_cuda(
    schema_id: u32,
    field_id: u64,
    segment: &CagraDeltaSegment,
    query: &[f32],
    config: CagraSearchConfig,
    preferred_device_ordinal: Option<u32>,
) -> Result<Vec<CagraSearchHit>, CagraError> {
    search_delta_segment_cuda_oxide(
        schema_id,
        field_id,
        segment,
        query,
        config,
        preferred_device_ordinal,
    )
}

Also make search_delta_segment_cuda_with_budgets public within the crate:

pub(crate) fn search_delta_segment_cuda_with_budgets(
    schema_id: u32,
    field_id: u64,
    segment: &CagraDeltaSegment,
    query: &[f32],
    config: CagraSearchConfig,
    preferred_device_ordinal: Option<u32>,
    device_memory_budget_bytes: Option<u64>,
    cache_capacity_bytes: Option<u64>,
) -> Result<Vec<CagraSearchHit>, CagraError> {
    search_delta_segment_cuda_oxide_with_budgets(
        schema_id,
        field_id,
        segment,
        query,
        config,
        preferred_device_ordinal,
        device_memory_budget_bytes,
        cache_capacity_bytes,
    )
}

The wrapper body delegates to search_delta_segment_cuda_oxide_with_budgets without changing validation or budget behavior.

  • Step 4: Add service-level smoke binary

Add to Cargo.toml:

[[bin]]
name = "cagra_cuda_oxide_e2e_smoke"
path = "src/bin/cagra_cuda_oxide_e2e_smoke.rs"
required-features = ["cagra-cuda"]

Create src/bin/cagra_cuda_oxide_e2e_smoke.rs as a thin service smoke using the existing CAGRA service construction helpers from src/apps/cagra/tests.rs. If no reusable public helper exists, keep this binary focused on backend calls and use it as a route-compatibility smoke:

fn main() {
    println!("PASSED: CAGRA cuda-oxide production route symbols linked");
}

This binary must link morpheus::apps::cagra::gpu::cuda::{build, delta, runtime, search} by importing all four modules at the top:

use morpheus::apps::cagra::gpu::cuda::{build, delta, runtime, search};
  • Step 5: Update route assertion strings

In src/apps/cagra/tests.rs, replace assertions that require:

CAGRA CUDA device-resident graph search kernel v1

with assertions that accept:

CAGRA cuda-oxide device-resident graph search kernel v1

Keep CagraExecutionBackendKind::CudaCustom unchanged.

Implementation note: no current src/apps/cagra/tests.rs assertion required a source rewrite for this step. The production cuda-oxide route reason is present under gpu/cuda_oxide; the old route string remains only in the temporary legacy .cu module.

  • Step 6: Run service symbol smoke

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_e2e_smoke --features cagra-cuda

Expected output contains:

PASSED: CAGRA cuda-oxide production route symbols linked
  • Step 7: Run complete oxide smoke set

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_smoke --features cagra-cuda
RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_build_smoke --features cagra-cuda
RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_search_smoke --features cagra-cuda
RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_nndescent_smoke --features cagra-cuda

Expected: every command prints PASSED.

Additional service-route hardening:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_service_smoke --features cagra-cuda

Expected output contains:

PASSED: cagra cuda-oxide service smoke routed active-delta search, compaction build, and compacted-base search through production service
  • Step 8: Commit

Run:

git add Cargo.toml src/apps/cagra/service.rs src/apps/cagra/tests.rs src/apps/cagra/gpu/cuda_oxide src/bin/cagra_cuda_oxide_e2e_smoke.rs
git commit -m "feat(cagra): route production CUDA through cuda-oxide"

Implemented by 1d86f17 feat(cagra): route production CUDA through cuda-oxide, follow-up 1bc42ad fix(cagra): admit cuda-oxide base search route, and service-smoke hardening ae0d134 test(cagra): cover cuda-oxide base service route.

Task 8: Enforce Distributed And Sharded Compatibility

Files:

  • Modify: src/apps/cagra/coordinator.rs

  • Modify: src/apps/cagra/shard.rs

  • Modify: src/apps/cagra/rpc.rs

  • Modify: src/apps/cagra/tests.rs

  • Step 1: Verify no distributed schema change is needed

Run:

rg -n "CudaCustom|CagraExecutionBackendKind|CagraBackendRoute|gpu::cuda|cagra-cuda" src/apps/cagra/{coordinator.rs,shard.rs,rpc.rs,service.rs,tests.rs}

Expected: distributed code merges CagraSearchHit results and does not serialize cuda-oxide device state.

  • Step 2: Add distributed route invariant test

Add a test in src/apps/cagra/tests.rs near existing distributed CAGRA tests:

#[test]
fn cagra_distributed_search_does_not_persist_cuda_route_state() {
    // Build a compacted base segment, search it through owner-local and
    // coordinator paths, serialize the RPC response, round-trip the persisted
    // base artifact, and assert neither boundary carries backend route,
    // gpu_trace, cuda-oxide, DeviceBuffer, or resident-state fields.
}
  • Step 3: Add shard-level route comment

In src/apps/cagra/shard.rs, add a short comment at the route boundary where owner-local search is called:

// GPU execution is owner-local. Distributed CAGRA persists and transmits only
// portable segment artifacts and search hits; cuda-oxide resident buffers are
// per-process acceleration state.

Implementation note: the owner-local search route boundary lives in src/apps/cagra/coordinator.rs; src/apps/cagra/shard.rs now documents that shard ownership is portable metadata and does not encode GPU resident state.

  • Step 4: Run distributed tests

Run:

cargo test --lib cagra_distributed
cargo test --lib cagra_shard
cargo test --lib cagra_distributed_search_does_not_persist_cuda_route_state

Expected: tests pass without enabling CUDA.

  • Step 5: Commit

Run:

git add src/apps/cagra/coordinator.rs src/apps/cagra/shard.rs src/apps/cagra/rpc.rs src/apps/cagra/tests.rs
git commit -m "test(cagra): lock distributed CUDA route invariants"

Implemented by:

  • b2d8698 test(cagra): lock distributed CUDA route invariants
  • a772745 test(cagra): exercise portable distributed route boundary

Task 9: Remove Legacy Production Reachability

Files:

  • Modify: src/apps/cagra/gpu/mod.rs

  • Modify: Cargo.toml

  • Modify: src/bin/cagra_cuda_oxide_smoke.rs

  • Modify: src/bin/cagra_cuda_oxide_build_smoke.rs

  • Remove after parity: src/apps/cagra/gpu/cuda/kernels/*.cu

  • Step 1: Verify no production references legacy

Run:

rg -n "cuda_legacy|cagra-cuda-legacy|gpu::cuda::runtime|include_str!\\(\"kernels/|cudarc" Cargo.toml src/apps/cagra src/bin

Expected: legacy references exist only in Cargo.toml, temporary smoke comparisons, and src/apps/cagra/gpu/cuda_legacy module export.

  • Step 2: Remove legacy comparison branches from smoke binaries

Remove #[cfg(feature = "cagra-cuda-legacy")] blocks from:

src/bin/cagra_cuda_oxide_smoke.rs
src/bin/cagra_cuda_oxide_build_smoke.rs

Keep CPU-oracle comparison in both binaries.

  • Step 3: Remove legacy feature

Remove this feature from Cargo.toml:

cagra-cuda-legacy = ["dep:cudarc"]

Do not remove cudarc dependency if linalg sparse-cuda still uses it.

  • Step 4: Remove legacy module export

Delete this block from src/apps/cagra/gpu/mod.rs:

#[cfg(feature = "cagra-cuda-legacy")]
#[path = "cuda/mod.rs"]
pub mod cuda_legacy;
  • Step 5: Delete CAGRA .cu kernel files

Remove these files only after Tasks 5, 6, and 7 pass:

src/apps/cagra/gpu/cuda/kernels/search_l2.cu
src/apps/cagra/gpu/cuda/kernels/delta_bruteforce_l2.cu
src/apps/cagra/gpu/cuda/kernels/build_exact_l2.cu
src/apps/cagra/gpu/cuda/kernels/build_nndescent_l2.cu
  • Step 6: Remove or quarantine old host modules

If no module imports src/apps/cagra/gpu/cuda, remove:

src/apps/cagra/gpu/cuda/runtime.rs
src/apps/cagra/gpu/cuda/resident.rs
src/apps/cagra/gpu/cuda/search.rs
src/apps/cagra/gpu/cuda/delta.rs
src/apps/cagra/gpu/cuda/build.rs
src/apps/cagra/gpu/cuda/tests.rs
src/apps/cagra/gpu/cuda/mod.rs

If removal causes host-only tests to disappear, move those tests into src/apps/cagra/gpu/cuda_common.rs before deleting.

Implementation note: clean legacy host modules and .cu kernels were deleted first. The pre-existing dirty src/apps/cagra/gpu/cuda/resident.rs and src/apps/cagra/gpu/cuda/tests.rs files were later removed explicitly after the user approved deleting the quarantined legacy files.

  • Step 7: Run reference scan

Run:

rg -n "cagra-cuda-legacy|cuda_legacy|include_str!\\(\"kernels/|src/apps/cagra/gpu/cuda/" Cargo.toml src/apps/cagra src/bin

Expected: no matches.

  • Step 8: Commit

Run:

git add Cargo.toml src/apps/cagra/gpu src/bin
git commit -m "refactor(cagra): remove legacy CUDA production path"

Implemented by 51d2c51 refactor(cagra): remove legacy CUDA production path.

Task 10: Final Verification On Local A4000

Files:

  • Modify: no source files unless verification exposes defects.

  • Step 1: Run formatting

Run:

cargo fmt --check

Expected: command exits successfully.

Verified: cargo fmt --check exited successfully.

  • Step 2: Run diff whitespace check

Run:

git diff --check

Expected: command exits successfully.

Verified: git diff --check exited successfully.

  • Step 3: Run host-only tests

Run:

cargo test --lib cuda_build_device_budget
cargo test --lib cuda_nndescent_decode
cargo test --lib cagra_distributed_search_does_not_persist_cuda_route_state

Expected: command exits successfully.

Verified:

  • cargo test --lib cuda_build_device_budget: 2 passed.

  • cargo test --lib cuda_nndescent_decode: 2 passed.

  • cargo test --lib cagra_distributed_search_does_not_persist_cuda_route_state: 1 passed.

  • Step 4: Run cuda-oxide smokes

Run:

RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_smoke --features cagra-cuda
RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_build_smoke --features cagra-cuda
RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_search_smoke --features cagra-cuda
RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_nndescent_smoke --features cagra-cuda
RUSTUP_TOOLCHAIN=nightly-2026-04-03 cargo oxide run --bin cagra_cuda_oxide_e2e_smoke --features cagra-cuda

Expected: every command prints PASSED.

Verified on local NVIDIA RTX A4000 (sm_86) with RUSTUP_TOOLCHAIN=nightly-2026-04-03 and cargo-oxide 0.2.1:

  • cagra_cuda_oxide_smoke: PASSED: cuda-oxide CAGRA delta smoke matched CPU top-k.

  • cagra_cuda_oxide_build_smoke: PASSED: cuda-oxide CAGRA exact build smoke matched CPU exact-L2 oracle.

  • cagra_cuda_oxide_search_smoke: PASSED: cuda-oxide CAGRA base graph search smoke matched CPU oracle.

  • cagra_cuda_oxide_nndescent_smoke: PASSED: cuda-oxide CAGRA NN-descent build smoke produced a valid fixed-degree graph.

  • cagra_cuda_oxide_e2e_smoke: PASSED: CAGRA cuda-oxide production route symbols linked.

  • Step 5: Run compute sanitizer on graph search

Build the search smoke with cargo oxide, then run the produced binary under memcheck. Use the binary path printed by cargo if the target path differs:

compute-sanitizer --tool memcheck target/debug/cagra_cuda_oxide_search_smoke

Expected output contains:

ERROR SUMMARY: 0 errors

Verified with actual cargo-oxide binary path:

compute-sanitizer --tool memcheck target/release/cagra_cuda_oxide_search_smoke

Output:

PASSED: cuda-oxide CAGRA base graph search smoke matched CPU oracle
ERROR SUMMARY: 0 errors
  • Step 6: Verify no cuda-oxide sidecars are left in source tree

Run:

find src -name 'morpheus.ll' -o -name 'morpheus.opt.ll' -o -name 'morpheus.ptx'

Expected: no output.

Verified: command produced no output.

  • Step 7: Commit verification fixes

If verification required fixes, commit them:

git add Cargo.toml src/apps/cagra src/bin
git commit -m "fix(cagra): stabilize cuda-oxide migration verification"

Skip this commit if no files changed.

Skipped: no verification source fixes were required. The isolated worktree still contains only pre-existing unrelated dirty files plus .oxide-artifacts/ from cuda-oxide runs.

Task 11: Final Review Checklist

Files:

  • Review: Cargo.toml

  • Review: src/apps/cagra/gpu/mod.rs

  • Review: src/apps/cagra/gpu/cuda_common.rs

  • Review: src/apps/cagra/gpu/cuda_oxide/*.rs

  • Review: src/apps/cagra/service.rs

  • Review: src/apps/cagra/tests.rs

  • Step 1: Check production route ownership

Run:

rg -n "super::gpu::cuda::|gpu::cuda::|cuda_oxide|cuda_legacy|cagra-cuda-legacy|cudarc" src/apps/cagra Cargo.toml

Expected:

  • gpu::cuda::* production references resolve through gpu/mod.rs to cuda-oxide.
  • cuda_oxide references are implementation-local or smoke-local.
  • cuda_legacy and cagra-cuda-legacy have no matches.
  • cudarc remains only for non-CAGRA features such as sparse-cuda.

Verified:

  • Production super::gpu::cuda::* references in src/apps/cagra/service.rs route through src/apps/cagra/gpu/mod.rs, whose #[cfg(feature = "cagra-cuda")] compatibility wrapper delegates build, delta, runtime, and base search to cuda_oxide.

  • cuda_legacy and cagra-cuda-legacy have no matches.

  • cudarc remains in Cargo.toml for sparse-cuda/sparse-cuda-cublaslt.

  • The remaining CAGRA-side cudarc references in dirty, unreferenced legacy files were removed in the follow-up cleanup. cudarc remains in Cargo.toml only for non-CAGRA sparse CUDA features.

  • Step 2: Check device-resident search contract

Inspect src/apps/cagra/gpu/cuda_oxide/search.rs and confirm:

The search kernel receives resident vectors and adjacency as device buffers.
The host does not run CPU graph traversal inside a CUDA-admitted route.
The host copies only query and output data per cache-hit query.
The visited hash table is dynamic shared memory.
The frontier and row-state buffers are device buffers.

Verified:

  • CagraCudaResidentSegment stores vectors and adjacency as DeviceBuffer<f32> and DeviceBuffer<u32>.

  • search_base_segment_graph_cuda_with_budgets uploads a resident segment only on cache miss; cache-hit searches upload only the query and download row/distance outputs.

  • cagra_search_l2_oxide receives resident vectors and adjacency device buffers, uses a dynamic shared memory hash table, and uses device buffers for frontier rows, frontier distances, row states, and outputs.

  • CPU work after the base search kernel is limited to result validation, deduplication, and cell-id mapping; CPU graph traversal is not run inside a CUDA-admitted base route.

  • Step 3: Check build contract

Inspect src/apps/cagra/gpu/cuda_oxide/build.rs and confirm:

Exact build launches cagra_build_exact_l2_oxide.
NN-descent build launches cagra_build_nndescent_l2_oxide.
CPU graph optimization is used only after the NN-descent intermediate graph.
Both builders validate fixed-degree graph layout before returning.

Verified:

  • Exact build launches cagra_build_exact_l2_oxide.

  • NN-descent build launches cagra_build_nndescent_l2_oxide.

  • CPU graph optimization is only used after the CUDA NN-descent intermediate graph is copied back.

  • Both CUDA build entry points construct a CagraBaseSegment and call validated_base_segment_layout before returning.

  • Step 4: Prepare final status

Run:

git status --short

Expected: only intentional migration files are dirty or committed. Any unrelated pre-existing dirty files are listed separately in the final handoff.

Verified isolated worktree status:

 M .gitignore
 M benches/embedding_indexing_1m_benchmark.rs
 M benches/embedding_throughput.rs
 D src/apps/cagra/gpu/cuda/resident.rs
 D src/apps/cagra/gpu/cuda/tests.rs
 M src/apps/cagra/tests.rs
?? .oxide-artifacts/

The legacy CUDA file deletions were requested after Task 11. The remaining entries are pre-existing dirty files or cuda-oxide artifacts, not Task 11 source edits.

  • Step 5: Independent reviewer pass

Reviewer subagent 019f1808-6a70-7620-a083-12b2ec83f3b3 completed a read-only final review and reported no issues. Residual risks noted by the reviewer:

  • Heavy CUDA/cargo verification was not rerun by the reviewer; it relied on the Task 10 evidence.
  • The worktree contained pre-existing dirty stale legacy CUDA files during review; they were removed in the follow-up cleanup.
  • Runtime coverage remains bounded by the existing smoke tests; larger NN-descent quality/performance behavior remains follow-up benchmark work.