cuda_graph_t discards CUDA error codes, obscuring capture failures

[ramakrishnap-nv]

Summary

cuopt::routing::detail::cuda_graph_t (cpp/src/routing/cuda_graph.cuh) discards the return code of every CUDA call it makes:

void start_capture(rmm::cuda_stream_view stream)
{
  cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal);
  capture_started = true;
}

void end_capture(rmm::cuda_stream_view stream)
{
  ...
  cudaStreamEndCapture(stream, &graph);
  ...
  cudaGraphExecUpdate(instance, graph, &errorNode, &updateResult);
  ...
  cudaGraphInstantiate(&instance, graph);
  cudaGraphDestroy(graph);
}

void launch_graph(rmm::cuda_stream_view stream) { cudaGraphLaunch(instance, stream); }

When a capture is invalidated mid-flight, cudaStreamEndCapture returns an error and sets graph to nullptr. That null graph is then passed to cudaGraphInstantiate, and instance is later launched, all unchecked.

Why this matters

This actively obstructed the diagnosis of #1748. The real failure was a single cudaErrorStreamCaptureUnsupported at one line inside a capture region. What CI showed instead was a cascade of downstream cudaErrorStreamCaptureInvalidated errors from unrelated call sites, because the first error was swallowed and execution continued on a null graph.

The failure was found in the end, but the logs pointed away from the cause rather than at it.

Prior attempt and why it was reverted

#1753 initially added RAFT_CUDA_TRY around these calls. That was scope creep on a CI hotfix and was reverted in favour of keeping that PR to the actual fix (see the review discussion there).

The attempt also surfaced two things worth carrying into any future change — both were bugs introduced by adding the checks naively, and both are traps for whoever does this properly:

1. capture_started must be cleared before the error propagates. RAFT_CUDA_TRY(cudaStreamEndCapture(...)) throws before the flag is reset, leaving a reused wrapper with stale capture state. Store the code, clear the flag, then throw.

2. A failed cudaGraphInstantiate leaks the captured graph. Throwing skips the cudaGraphDestroy(graph) at the end of end_capture. Destroy before propagating, or use a scope guard.

cpp/src/utilities/manual_cuda_graph.cuh already handles both correctly and is the reference for the ordering.

Open questions for whoever picks this up

• cudaGraphExecUpdate. Its failure is a normal, expected path — the code falls back to re-instantiation. But the current code cannot distinguish "update failed, re-instantiate" from a genuine error, since it only inspects updateResult and never the API return code. cudaErrorGraphExecUpdateFailure should feed the existing fallback; anything else is probably a real error.
• Should launch_graph verify graph_created? Raised in review on fix: avoid uncapturable RMM copies inside routing CUDA graph captures #1753: if instantiation failed, instance may be unset or destroyed.
• Testing. There is currently no coverage of these paths, and exercising a failed cudaGraphInstantiate needs CUDA fault injection, which cpp/tests does not have. Note also that cpp/tests/routing/local_search_cand_test.cu is not registered in cpp/tests/routing/CMakeLists.txt.

Related

• [BUG] All CUDA 13 test jobs fail: RMM allocation during CUDA graph capture (cudaErrorStreamCaptureUnsupported) #1748 — the CUDA 13 capture failures this obscured
• fix: avoid uncapturable RMM copies inside routing CUDA graph captures #1753 — the fix, where the error checks were reverted out
• [BUG] cudaMemcpyBatchAsync path from #2511 breaks CUDA graph capture on CUDA 13 (silently drops the copy) rapidsai/rmm#2518 — the upstream cause