[BUG] test_batch_solve_varying_sizes: block_copy destination too small on CUDA 13.3.0 (blocks all PR wheel CI)

Summary

test_batch_solve_varying_sizes fails on CUDA 13.3.0 with a device-side assert in block_copy: the destination span is smaller than the number of elements being copied.

cpp/src/utilities/cuda_helpers.cuh:147: void cuopt::block_copy(
  raft::span<T,...>, raft::span<T,...>, unsigned long) [with T = cuopt::routing::detail::NodeInfo<int>]:
  block: [0,0,0], thread: [192,0,0]
  Assertion `dst.size() >= size && "block_copy::dst does not have the sufficient size"` failed.

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  fill_n: failed to synchronize: cudaErrorAssert: device-side assert triggered
[gw2] node down: Not properly terminated

The assert fires across hundreds of threads and aborts the process, taking the pytest-xdist worker down with it.

Observed in

Reproduces on any PR that runs the wheel tests. Two examples on unrelated branches:

#1753 job 96487821122 — CUDA graph capture fix
#1726 job 96126562276 — Mod-2 cut separator

Neither touches routing batch solve, so this is not caused by either.

It is specific to CUDA 13.3.0

From a single PR run, wheel-tests-cuopt:

CUDA	result
13.3.0	3 / 3 fail
13.0.3	2 / 2 pass
12.9.2	pass

13.0.3 passes, so this is not "CUDA 13" generally — something about the 13.3 toolkit specifically.

Why only wheel jobs, and why nightly is green

Assertions are enabled only for pull-request wheel builds:

# ci/build_wheel_libcuopt.sh
if [ "$RAPIDS_BUILD_TYPE" = "pull-request" ]; then
    export SKBUILD_CMAKE_ARGS="${SKBUILD_CMAKE_ARGS};-DDEFINE_ASSERT=True"

cuopt_assert compiles to nothing without ASSERT_MODE (cpp/src/utilities/macros.cuh:15-22). Conda builds never set DEFINE_ASSERT, and nightly wheel builds take the release branch.

Warning

This means the undersized copy is almost certainly also happening in release and conda builds — silently. There the assert is compiled out and block_copy writes past the end of the destination span. The failing PR jobs are the only place we currently see it; a green nightly does not mean the code path is correct.

That should be confirmed or ruled out (e.g. compute-sanitizer on this test with a release build) before treating this as test-only.

Call path

route.cuh:541  requests().copy_fixed_route_data(orig_route.requests(), start_idx, end_idx, write_start)
  pdp_route.cuh:93   block_copy(node_info.subspan(write_start),
                                orig_route.node_info.subspan(from_idx),
                                to_idx - from_idx)
    cuda_helpers.cuh:147  cuopt_assert(dst.size() >= size, ...)   <-- fails

Suspected cause

The test batch-solves TSPs of interleaved varying sizes:

tsp_sizes = [5, 8, 10, 6, 7, 9, 12, 15, 11, 4, 13, 14, 8, 6, 10, 9, 7, 11, 5, 12]
solutions = routing.BatchSolve(data_models, settings)

A route/solution buffer sized for one problem being reused for a larger problem later in the batch would produce exactly this: write_start + size exceeding the destination span. Not yet confirmed — the sizing logic in the batch path needs a proper look.

Unclear why CUDA 13.3 specifically surfaces it; that may just be scheduling/timing exposing a latent bug rather than a 13.3 regression.

Impact on CI

Every PR that runs wheel-tests-cuopt on 13.3.0 is red because of this, regardless of what the PR changes. It is currently the only failure on #1753.

Suggested next steps

Confirm whether the out-of-bounds write also occurs in release builds (compute-sanitizer)
Find the buffer whose size is computed from the wrong problem in the batch path
Determine why 13.3.0 exposes it and 13.0.3 does not
Temporarily skip the test to unblock CI (note: xfail will not work — the device assert aborts the process, so pytest cannot catch it)

fix: read max_active_nodes with raft::copy instead of device_scalar::…

test: skip test_batch_solve_varying_sizes pending NVIDIA#1756

Update: test skipped, and this is a regression rather than a latent bug

test_batch_solve_varying_sizes is now skipped as of 402005c1 on #1753, so PR CI is no longer blocked by it. This issue now also tracks re-enabling that test — the skip must be removed as part of the fix.

skip rather than xfail is deliberate: the device-side assert aborts the process, so pytest cannot catch it, and run_cuopt_pytests.sh runs with --max-worker-restart=0.

New evidence: it started in the rmm#2511 window

wheel-tests-cuopt / 13.3.0 (all three jobs, assertions enabled since these are PR builds) was green on PRs merged immediately before rapidsai/rmm#2511 landed on 2026-08-18:

PR	13.3.0 wheel jobs
#1723	pass
#1713	pass
#1714	pass

So this is not a long-standing latent bug that CUDA 13.3 happens to expose. Something changed in that window. rmm#2511 rerouted rmm::detail::memcpy_async to cudaMemcpyBatchAsync on CUDA 13, which is the same change behind #1748 / rapidsai/rmm#2518.

Hypothesis tested and refuted

device_scalar::value() performs its device-to-host read through rmm::detail::memcpy_async (rmm/device_uvector.hpp:311), and compute_max_active uses it to read max_active_nodes, which sizes the route buffers that later overflow:

// solution.cu:713-719
compute_max_active_kernel<<<1, TPB, 0, sol_handle->get_stream()>>>(view());
max_active_nodes = max_active_nodes_for_all_routes.value(sol_handle->get_stream());

Replacing that read with raft::copy (plain cudaMemcpyAsync) — commit 5a93131a on #1753 — did not fix it. CI ran that exact commit and the three 13.3.0 jobs failed identically: same test, same 256 block_copy::dst asserts, 1 failed, 69 passed.

So the stale-D2H-read theory is wrong, at least for this call site. Recording it so nobody repeats it.

Suggested next step

Pin librmm to a pre-#2511 build (e.g. 26.10.00a24) on a scratch branch and re-run the 13.3.0 wheel jobs.

If they go green, rmm is confirmed as the trigger and this becomes evidence for [BUG] cudaMemcpyBatchAsync path from #2511 breaks CUDA graph capture on CUDA 13 (silently drops the copy) rapidsai/rmm#2518.
If they still fail, the Aug-18 correlation is a coincidence and this is a cuOpt bug that 13.3 exposes — at which point the productive move is to instrument the assert site and print dst.size(), write_start, size and the batch index from CI, rather than guessing at further substitutions.

That ordering matters because one substitution-based guess has already failed.

Reminder on severity

Assertions are compiled out in conda and nightly builds (ci/build_wheel_libcuopt.sh only sets -DDEFINE_ASSERT=True for RAPIDS_BUILD_TYPE = pull-request). The undersized copy is very likely still happening there, silently writing past the end of the span. With the test now skipped, there is no signal at all — so this should not be treated as a CI-only annoyance.

Not reproducible on a CUDA 13.0 driver

Local repro attempts were inconclusive because the driver caps the effective runtime: on driver 580 (CUDA 13.0) with a 13.3 toolkit and an ASSERT_MODE build, the test passes deterministically. That is the same configuration as the passing 13.0.3 CI jobs. Reproducing needs driver 595+ (what latest-driver runners use), not just a newer toolkit.

diag: print sizes and launch geometry on block_copy overflow

diag: add route-level and host->device n_nodes checks for NVIDIA#1756

fix: keep host sources alive for async route-view copies (NVIDIA#1756)

refactor: publish route views through a single path

fix: keep host sources alive for async route-view copies (NVIDIA#1756) (

Navigation Menu

Repository navigation

[BUG] test_batch_solve_varying_sizes: block_copy destination too small on CUDA 13.3.0 (blocks all PR wheel CI) #1756

Summary

Observed in

It is specific to CUDA 13.3.0

Why only wheel jobs, and why nightly is green

Call path

Suspected cause

Impact on CI

Suggested next steps

Update: test skipped, and this is a regression rather than a latent bug

New evidence: it started in the rmm#2511 window

Hypothesis tested and refuted

Suggested next step

Reminder on severity

Not reproducible on a CUDA 13.0 driver

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

[BUG] test_batch_solve_varying_sizes: block_copy destination too small on CUDA 13.3.0 (blocks all PR wheel CI) #1756

Description

Summary

Observed in

It is specific to CUDA 13.3.0

Why only wheel jobs, and why nightly is green

Call path

Suspected cause

Impact on CI

Suggested next steps

Activity

ramakrishnap-nv commented on Aug 20, 2026

Update: test skipped, and this is a regression rather than a latent bug

New evidence: it started in the rmm#2511 window

Hypothesis tested and refuted

Suggested next step

Reminder on severity

Not reproducible on a CUDA 13.0 driver

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions