bugfix: handle cuda graph errors

[neuriv]

Description

Check CUDA graph return values and keep capture and executable-graph state consistent on failure. Release temporary graphs with the existing scope guard. Preserve successful updates and clear expected update errors before rebuilding an executable graph.

Closes #1758

Testing

ubuntu 22.04, cuda 12.9.86, 1xa100.

cmake --build "$build" --target ROUTING_UNIT_TEST --parallel 20

filter='objective_function.total_time:vehicle_order_match.one_order_per_vehicle:route_constraints.vanilla_pdp:batch_tsp.varying_sizes'
"$build/tests/routing/ROUTING_UNIT_TEST" \
  --gtest_filter="RoutingCudaGraph.*:$filter"

compute-sanitizer --tool memcheck --report-api-errors no \
  --leak-check full --error-exitcode 99 \
  "$build/tests/routing/ROUTING_UNIT_TEST" \
  --rmm_mode=cuda --gtest_filter='RoutingCudaGraph.*'

pre-commit run --all-files

test passed. compute sanitizer shows 0 issues as well.
The four routing integration tests passed on baseline.

unit test

#!/usr/bin/env bash
set -euo pipefail
repo=$(cd "${1:?provide the cuopt checkout}" && pwd)
gtest=$(cd "${2:?provide the googletest checkout}" && pwd)
: "${CONDA_PREFIX:?activate the CUDA development environment}"
test_dir=$(mktemp -d)
trap 'rm -rf "$test_dir"' EXIT
cat > "$test_dir/cuda_graph.cu" <<'CPP'
/* clang-format off */
/*
 * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: Apache-2.0
 */
/* clang-format on */

#include <routing/cuda_graph.cuh>

#include <raft/core/handle.hpp>
#include <rmm/device_uvector.hpp>

#include <gtest/gtest.h>

namespace cuopt {
namespace routing {
namespace test {

class RoutingCudaGraph : public ::testing::Test {
 protected:
  void TearDown() override
  {
    if (graph.graph_created) { EXPECT_EQ(cudaSuccess, cudaGraphExecDestroy(graph.instance)); }
    cudaGetLastError();
  }

  raft::handle_t handle;
  detail::cuda_graph_t graph{};
};

__global__ void write_graph_value(int* output, int value) { *output = value; }

TEST_F(RoutingCudaGraph, BeginCaptureFailure)
{
  EXPECT_THROW(graph.start_capture(cuda::stream_ref{cudaStreamLegacy}), raft::cuda_error);
  EXPECT_FALSE(graph.capture_started);
}

TEST_F(RoutingCudaGraph, InvalidatedCaptureCanRetry)
{
  auto stream = handle.get_stream();
  graph.start_capture(stream);
  EXPECT_EQ(cudaErrorStreamCaptureUnsupported, cudaStreamSynchronize(stream.get()));
  EXPECT_THROW(graph.end_capture(stream), raft::cuda_error);
  EXPECT_FALSE(graph.capture_started);
  EXPECT_FALSE(graph.graph_created);
  EXPECT_EQ(graph.graph, nullptr);
  cudaGetLastError();

  graph.start_capture(stream);
  graph.end_capture(stream);
  EXPECT_TRUE(graph.graph_created);
  EXPECT_FALSE(graph.capture_started);
  graph.launch_graph(stream);
  handle.sync_stream();
}

TEST_F(RoutingCudaGraph, CaptureRequired)
{
  EXPECT_THROW(graph.end_capture(handle.get_stream()), cuopt::logic_error);
  EXPECT_THROW(graph.launch_graph(handle.get_stream()), cuopt::logic_error);
}

TEST_F(RoutingCudaGraph, UpdateAndReplacement)
{
  auto stream = handle.get_stream();
  rmm::device_uvector<int> output(1, stream);
  for (int value = 1; value <= 3; ++value) {
    graph.start_capture(stream);
    write_graph_value<<<1, 1, 0, stream.get()>>>(output.data(), value);
    if (value == 3) { write_graph_value<<<1, 1, 0, stream.get()>>>(output.data(), value); }
    graph.end_capture(stream);
    ASSERT_TRUE(graph.graph_created);
    EXPECT_FALSE(graph.capture_started);
    EXPECT_EQ(cudaSuccess, cudaGetLastError());
    if (value == 2) { EXPECT_EQ(cudaGraphExecUpdateSuccess, graph.updateResult); }
    if (value == 3) { EXPECT_NE(cudaGraphExecUpdateSuccess, graph.updateResult); }

    graph.launch_graph(stream);
    int result = 0;
    RAFT_CUDA_TRY(cudaMemcpyAsync(
      &result, output.data(), sizeof(result), cudaMemcpyDeviceToHost, stream.get()));
    handle.sync_stream();
    EXPECT_EQ(value, result);
  }
}

}  // namespace test
}  // namespace routing
}  // namespace cuopt
CPP
nvcc -std=c++20 -arch="sm_${CUDA_ARCH:-80}" -Xcompiler=-pthread \
  --expt-relaxed-constexpr \
  -I"$repo/cpp/include" -I"$repo/cpp/src" \
  -I"$CONDA_PREFIX/include" -I"$CONDA_PREFIX/include/rapids" \
  -I"$gtest/googletest/include" -I"$gtest/googletest" \
  "$test_dir/cuda_graph.cu" "$gtest/googletest/src/gtest-all.cc" \
  "$gtest/googletest/src/gtest_main.cc" \
  -L"$CONDA_PREFIX/lib" -lrmm -lrapids_logger \
  -Xlinker -rpath -Xlinker "$CONDA_PREFIX/lib" -o "$test_dir/unit-test"
"$test_dir/unit-test"
compute-sanitizer --tool memcheck --report-api-errors no \
  --leak-check full --error-exitcode 99 "$test_dir/unit-test"

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Repository: NVIDIA/cuopt/.coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 6409ceba-b9be-4478-8d14-7e6ff656ba38

📥 Commits

Reviewing files that changed from the base of the PR and between 1d591b2 and b538c3d.

📒 Files selected for processing (1)

• cpp/src/routing/cuda_graph.cuh

Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.

---

📝 Walkthrough

Walkthrough

Changes

CUDA graph error handling

Layer / File(s)	Summary
Capture completion and graph instantiation cpp/src/routing/cuda_graph.cuh	CUDA capture and graph operations now use RAFT error handling. end_capture validates capture state, handles expected update failures, cleans up captured graphs, and re-instantiates outdated executable graphs.
Validated graph launch cpp/src/routing/cuda_graph.cuh	launch_graph now checks that graph instantiation completed before launching and reports CUDA launch errors.

Priority: ➖ Normal

Estimated code review effort: 2 (Simple) | ~10 minutes

Change: Bug fix · Severity of issue fixed: Medium

Suggested reviewers: bdice

Merge Risk: ⚪ Minimal · up to b538c

The changed graph lifecycle preserves valid launch state and propagates unexpected CUDA failures. No actionable merge risk remains.

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Linked Issues check	⚠️ Warning	Issue #1758 coding requirements are partly implemented. cpp/src/routing/cuda_graph.cuh checks capture, update, instantiation, destruction, and launch results. It clears capture_started before prop…	Add and register automated tests for the required failure and recovery paths in Issue #1758. The tests must verify error propagation, capture_started reset after end-capture failure, retry behavior, expected versus unexpected update error…

✅ Passed checks (4 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly and concisely describes the main change: handling CUDA graph errors.
Description check	✅ Passed	The description directly explains the CUDA graph error handling changes and includes relevant testing details.
Out of Scope Changes check	✅ Passed	The whole-PR diff changes only cpp/src/routing/cuda_graph.cuh. The include changes, CUDA error checks, state validation, and scope guard directly support Issue #1758. No unrelated production change …
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check. Docstring coverage is scoped to functions touched by this diff. Analyzed 0 functions across 0…

Full details: Linked Issues check

Explanation

Issue #1758 coding requirements are partly implemented. cpp/src/routing/cuda_graph.cuh checks capture, update, instantiation, destruction, and launch results. It clears capture_started before propagating end-capture errors. It uses a scope guard to destroy temporary graphs. It treats cudaErrorGraphExecUpdateFailure as the re-instantiation path and rejects launches before instantiation. However, the whole-PR diff contains no test changes. It does not add or register automated coverage for capture failures, retry recovery, required capture state, graph updates, and graph replacement.

Resolution

Add and register automated tests for the required failure and recovery paths in Issue #1758. The tests must verify error propagation, capture_started reset after end-capture failure, retry behavior, expected versus unexpected update errors, launch rejection before instantiation, graph cleanup after instantiation failure, graph updates, and graph replacement.

• Fix all pre-merge checks with AI

✨ Finishing Touches 💡 1

🛠️ Fix failing CI checks 💡

• Commit to this branch
• Create a new PR

🧪 Generate unit tests (beta)

• Create a new PR

---

_{Comment @coderabbitai help to get the list of available commands.}