[BUG] Multi-dimension capacity data corrupted when routes are resized mid-search (flaky `Cost should improve!` / `Cost mismatch` SIGABRT in ROUTING_TEST)

[mlubin]

Symptom

conda-cpp-tests (assert-enabled builds) occasionally crash with SIGABRT in ROUTING_TEST, e.g. this PR CI run on level0_retail/retail_float_test_t.CVRPTW_Retail/4:

vrp_execute.cu:466: execute_vrp_moves(...) [f_t = float, REQUEST = VRP]:
Assertion `cost_before - cost_after > EPSILON && "Cost should improve!"' failed.

The sibling assertion ("Cost mismatch on vrp costs!", vrp_execute.cu:467) fires for the same underlying reason. The failure is rare (~1% of runs of that single test locally) and timing-dependent, so it presents as a CI flake. In release builds (no asserts) the same bug does not crash — it silently corrupts the solution's capacity data mid-search.

Root cause

capacity_route_t stores its per-node arrays row-major with one row per capacity dimension:

• element access: demand[dim * stride + idx] (cpp/src/routing/route/capacity_route.cuh, get_node/set_node)
• stride is derived from the current allocation: view() sets stride = demand.size() / n_capacity_dimensions (capacity_route.cuh:202)

capacity_route_t::resize() (capacity_route.cuh:56–62) grows these buffers with a plain rmm::device_uvector::resize. When that reallocates, it preserves only the linear prefix of the old buffer; when the new size still fits within existing capacity, RMM moves nothing at all — but the stride derived from size() still changes, which shifts every row boundary either way. After growing from stride_old to stride_new:

• row 0 (first capacity dimension) still starts at offset 0 → intact;
• row i ≥ 1 was written at i * stride_old but is now read at i * stride_new → garbage for every capacity dimension beyond the first.

Nothing re-populates the data afterward: solution_t::resize_routes (solution.cu:373) resizes and returns. Callers rely on resize preserving data — add_route/add_routes grow all pre-existing routes without re-filling them, and check_routes_can_insert_and_get_sh_size is documented as "call before each kernel that changes route sizes", with the subsequent kernel immediately reading the existing route data.

Structurally, resize is the only member of capacity_route_t that treats the buffers as flat: get_node/set_node and even the partial-copy helpers are all stride-aware (the copy helpers re-stride between source and destination routes, capacity_route.cuh:112–147). Its body is the same boilerplate as the single-row dimensions (time_route.cuh:69, distance_route.cuh:50) with n_capacity_dimensions * multiplied into the allocation size only — which preserves byte count but not layout. The capacity dimension is the only one with a row-major multi-row layout, so every other dimension survives resizes untouched.

Trigger chain (why it's flaky)

1. VRP local search selects a 2-opt* move that grows a route substantially (captured instance: a tail swap growing a route from 12 to 38 nodes).
2. execute_vrp_moves → check_routes_can_insert_and_get_sh_size (solution.cu:364) → resize_routes(alignTo(max_active + added, base_route_size)), which resizes every route whose buffer is smaller than the new size. This only happens when a move outgrows the current aligned allocation — rare, and dependent on the (time-limited, randomized) search trajectory.
3. The execute kernel then reads the scrambled row-1 demands, computes garbage capacity excess, and the recomputed cost explodes; whichever of the two debug assertions trips first aborts the process.

Only problems with ≥ 2 capacity dimensions are affected. In the retail gtest suite that's parametrizations /1 (multi-capacity) and /4 (multi-capacity + vehicle time windows) — CI failures observed on /4.

Evidence

Reproduced locally on an RTX PRO 6000 Blackwell, CUDA 13.3, ./build.sh libcuopt -a, looping ROUTING_TEST --gtest_filter=level0_retail/retail_float_test_t.CVRPTW_Retail/4: 2 hits in ~101 runs (one per assert flavor). With diagnostic instrumentation added in execute_vrp_moves (which perturbs timing), the failure became near-deterministic (3 for 3). A captured failing execution shows:

VRP_EXEC_MARGIN realized=-2.67e+06 promised=+1.84 n_moves=1 cost_before=810.9
                excess_before=9.88e+01 excess_after=2.67e+06
VRP_EXEC_MOVE   type=133 (2-opt* w/ vehicle swap) r1=2 r2=4 fs1=1 fs2=1
route 2: pre_inf CAP=0        → post_inf CAP=1.083e+09   (grew 12 → 38 nodes)
route 4: pre_inf CAP=630      → post_inf CAP=1.118e+09

• The explosion is only in the CAP dimension (~2³⁰-scale garbage — the demand spans are i_t), on both routes (both were resized), while distance/time stay sane and the post-move node sequences exactly match what the evaluator scored.
• In runs where nothing fires, promised vs. realized deltas agree to 0.0e+00–1e-13 — the delta evaluation machinery itself is exact; the divergence appears only when a resize intervenes.
• One captured run had cost_before = 2.4e+07 before the failing move — i.e., the solution had already been silently poisoned by an earlier resize, which is the release-build failure mode.

Proposed fix

Make capacity_route_t::resize stride-aware: re-place each dimension's row at its new offset instead of flat-resizing (a cudaMemcpy2DAsync with spitch = stride_old, dpitch = stride_new does this in one call per array):

void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
{
  i_t n_dims = dim_info.n_capacity_dimensions;
  if (n_dims == 0) { return; }
  auto resize_strided = [&](rmm::device_uvector<i_t>& vec) {
    i_t old_stride = vec.size() / n_dims;
    i_t new_stride = max_nodes_per_route;
    if (old_stride == new_stride) { return; }
    rmm::device_uvector<i_t> new_vec(n_dims * new_stride, stream);
    if (old_stride > 0) {
      RAFT_CUDA_TRY(cudaMemcpy2DAsync(new_vec.data(), new_stride * sizeof(i_t),
                                      vec.data(), old_stride * sizeof(i_t),
                                      std::min(old_stride, new_stride) * sizeof(i_t),
                                      n_dims, cudaMemcpyDeviceToDevice, stream.value()));
    }
    vec = std::move(new_vec);
  };
  resize_strided(demand);
  resize_strided(gathered);
  resize_strided(max_to_node);
  resize_strided(max_after);
}

Verified locally (assert build, same GPU/CUDA as the failing CI leg):

Check	Pre-fix	Post-fix
CVRPTW_Retail/4 loop (instrumented build)	3/3 aborts	0/50 aborts
Full retail suite (19 params) ×3	—	all pass
ROUTING_UNIT_TEST	—	57/57 pass
Promised-vs-realized delta noise (675 near-boundary samples)	up to 2.2e+11	≤ 1e-13