[0z5a]

Description

cost_boundary_initialization.pdp_set_route_data_resets_cost_boundaries copied the pickup and delivery indices with raft::copy from stack locals declared inside the if constexpr (is_pdp) block, and the block ended before anything synchronized the stream. raft::copy is cudaMemcpyAsync with a pageable source, so on platforms where the driver reads the source when the stream reaches the copy, the device received stale stack contents and set_route_data read out of bounds (cudaErrorIllegalAddress). The sticky error then failed set_shmem_of_kernel.zero_request and normal_request. Only the PDP variant has these copies; the VRP variant passed in the same failing gb300 runs (10-25, 11-01). This is the same class of bug as #1759, #2200 and #2202.

Changes:

• cost_boundary_initialization.cu: build d_pickups/d_deliveries with cuopt::device_copy from function-scope vectors, as the test already does for cost_matrix and order_locations.
• routing/generator/generator.cu: generate_coordinates copied the depot coordinates from locals and returned without a sync; set them with thrust::fill_n instead.
• mip_heuristics/problem/problem.cu: csr_to_csc_transpose set the last CSC offset from a local the same way; same change.

The other raft::copy(..., &x, 1, ...) sites copy from members or parameters that outlive a later sync, so they are unchanged.

Testing: the three changed translation units compile with nvcc 13.3.1 against the 26.12.00a58 headers, and clang-format 20.1.8 leaves them unchanged. The failure does not reproduce on x86: on a B300 (sm_103, same sm_100f device code as GB300) with 26.12.00a58 the two test groups passed 200/200 before this change, and a small probe there shows pageable cudaMemcpyAsync reads its source at call time. So the gb300 nightly is the real check.

Issue

Refs #2212 (kept open until a gb300 nightly passes after merge)

Checklist

• I am familiar with the Contributing Guidelines.
• Testing
  • New or existing tests cover these changes
  • Added tests
  • Created an issue to follow-up
  • NA
• Documentation
  • The documentation is up to date with these changes
  • Added new documentation
  • NA

[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.

[0z5a]

Follow-up on the mechanism. I ran a small probe that queues a ~200 ms spin kernel, issues cudaMemcpyAsync(d, &src, 4, cudaMemcpyDefault, stream) from pageable memory, overwrites src right after the call, and checks what the device received. A pinned source is the positive control.

Machine	Driver	cudart	pageable stack / heap	pinned
B300, x86_64	610.57.04	13.3	read at call time	deferred
GH200, aarch64 (pageableMemoryAccessUsesHostPageTables=1)	595.58.03	13.2, 13.3, 13.4	read at call time	deferred

So neither machine defers the pageable read, which matches the gh200 nightly passing and the 50/50 passes I saw on GH200 earlier. The remaining difference from the failing runner is GB300 itself and/or its latest-driver, which I have no access to, so I can't confirm the deferred read there. The test does let the source go out of scope before the copy is known to have run, which CUDA only guarantees to be safe for pageable memory where the copy happens to be synchronous, so the change is correct either way; whether it is the cause of #2212 will show in the gb300 nightly.

probe

// Does cudaMemcpyAsync(H2D, pageable source) read the source at call time or when the stream reaches it?
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
__global__ void spin(long long cycles) { long long t = clock64(); while (clock64() - t < cycles) {} }
#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { printf("%s: %s\n", #x, cudaGetErrorString(e)); exit(1); } } while (0)
static int run(const char* name, int* src, int* d, cudaStream_t s, long long cyc) {
  *src = 111;
  spin<<<1, 1, 0, s>>>(cyc);                                   // stream busy for ~cyc cycles
  CK(cudaMemcpyAsync(d, src, sizeof(int), cudaMemcpyDefault, s));
  *src = 999;                                                  // source "dies" right after the call
  int h = 0;
  CK(cudaStreamSynchronize(s));
  CK(cudaMemcpy(&h, d, sizeof(int), cudaMemcpyDeviceToHost));
  printf("  %-18s device got %d -> %s\n", name, h, h == 999 ? "DEFERRED read (source read after call returned)" : "read at call time");
  return h == 999;
}
int main() {
  int dev = 0, drv = 0, rt = 0, pma = 0, pmaHPT = 0, clk = 0;
  CK(cudaSetDevice(dev)); cudaDriverGetVersion(&drv); cudaRuntimeGetVersion(&rt);
  cudaDeviceGetAttribute(&pma, cudaDevAttrPageableMemoryAccess, dev);
  cudaDeviceGetAttribute(&pmaHPT, cudaDevAttrPageableMemoryAccessUsesHostPageTables, dev);
  cudaDeviceGetAttribute(&clk, cudaDevAttrClockRate, dev);
  cudaDeviceProp p; cudaGetDeviceProperties(&p, dev);
  printf("%s sm_%d%d driver=%d runtime=%d pageableMemoryAccess=%d usesHostPageTables=%d\n", p.name, p.major, p.minor, drv, rt, pma, pmaHPT);
  long long cyc = (long long)clk * 1000LL / 5;                 // ~200 ms
  int* d; CK(cudaMalloc(&d, sizeof(int)));
  cudaStream_t s; CK(cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking));
  int stack_local; run("pageable stack", &stack_local, d, s, cyc);
  int* heap = (int*)malloc(sizeof(int)); run("pageable heap", heap, d, s, cyc);
  int* pinned; CK(cudaMallocHost(&pinned, sizeof(int))); run("pinned (control)", pinned, d, s, cyc);
  return 0;
}

[ramakrishnap-nv]

The diagnosis fits: the PDP-only raft::copy calls read stack locals that were gone by the time the stream ran, on a platform where the driver reads pageable sources late. I checked the other raft::copy(..., &x, 1, ...) sources and the ones this PR leaves alone are safe (members, a by-value parameter that is synced before return, or locals followed by a sync). Two small comments inline. The failure only shows on gb300, so please keep #2212 open until a gb300 nightly passes cleanly after merge (the PR body's Closes #2212 will close it on merge). The PR also needs the bug and non-breaking labels and a milestone.

[ramakrishnap-nv]

The VRP variant now also builds pickups/deliveries and calls device_copy on empty vectors. That is harmless, but it is work the VRP test did not do before. Optional: build and copy them only under if constexpr (is_pdp), or keep as is. If you keep it, please confirm device_copy on an empty vector does not allocate or launch anything that could change what the VRP test exercises.

[0z5a]

Done in eba2142: the host vectors stay at function scope (they have to outlive the copies), d_pickups/d_deliveries are device_uvector<int>(0, stream) as before the PR, and only the PDP branch fills them with device_copy. So the VRP variant does the same device work as on main: two empty device_uvectors, no copies.

[ramakrishnap-nv]

handle.get_thrust_policy() should run on handle.get_stream(), so ordering against the surrounding code is the same as the old raft::copy. Please confirm the returned v_x_pos/v_y_pos are only consumed on that stream (or after a sync), since the old code had the same ordering requirement.

[0z5a]

Confirmed. raft::resource::get_thrust_policy builds its rmm::exec_policy_nosync from get_cuda_stream(res), the stream handle.get_stream() returns, so fill_n is enqueued on the same stream as the old raft::copy. The callers of generate_coordinates are generate_dataset, generate_fleet_info and generate_time_windows in generator.cu, plus l0_vehicle_types_test.cu; each passes the coordinates to generate_matrices with the same handle, whose kernel reads them on handle.get_stream(). No caller reads them on another stream.

[ramakrishnap-nv]

Same as above for handle_ptr->get_thrust_policy(): please confirm it runs on the stream used for the surrounding scatter, so the last offset is written in order. total_nnz is passed by value into the kernel, so the lifetime problem is gone.

[0z5a]

Confirmed. csr_to_csc_transpose sets stream = handle_ptr->get_stream() at the top and already runs its thrust::fill, for_each and exclusive_scan through handle_ptr->get_thrust_policy(); the scatter kernel then launches on stream.get(). So the fill_n is ordered after the scan and before the scatter, on one stream.

[coderabbitai]

No actionable comments were generated in the recent review.

Walkthrough

The PR replaces host-to-device scalar copies with Thrust device fills for CSC offsets and depot coordinates. It also changes PDP test setup to keep host vectors alive during asynchronous device copies.

Changes: Device initialization

Layer / File(s)	Summary
Device-side value initialization cpp/src/mip_heuristics/problem/problem.cu, cpp/src/routing/generator/generator.cu	The CSC terminal offset and depot coordinates are initialized with thrust::fill_n instead of host-to-device scalar copies.
PDP test copy sources cpp/tests/routing/unit_tests/cost_boundary_initialization.cu	For PDP, the test copies pickup and delivery values from function-scope host vectors to device buffers. VRP skips this setup.

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

Change: Bug fix · Severity of issue fixed: Medium

Possibly related PRs

• NVIDIA/cuopt#2202: This PR also addresses asynchronous initialization and host-source lifetime hazards.

Suggested labels: bug, non-breaking

Merge Risk: Minimal · up to eba21

The changes preserve asynchronous copy source lifetimes and maintain stream ordering.

Pre-merge checks: 3 · 1 · 1

Failed checks: 1 warning · 1 inconclusive

Check name	Status	Explanation	Resolution
Docstring Coverage		Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff with cyclomatic complexity of at least 2. Documentat…	Write docstrings for the functions missing them to satisfy the coverage threshold.
Linked Issues check		Issue #2212 requires that cost_boundary_initialization.pdp_set_route_data_resets_cost_boundaries pass on gb300. The only coding change for this test is in cpp/tests/routing/unit_tests/cost_boundary_in…	Provide the gb300 CI result for cost_boundary_initialization.pdp_set_route_data_resets_cost_boundaries, set_shmem_of_kernel.zero_request, and set_shmem_of_kernel.normal_request on the reviewed head (eba2142). A passing result on gb300 resol…

Passed checks: 3 passed

Check name	Status	Explanation
Title check		The title names the main change: keeping host sources alive for async copies. The changeset replaces stack-local host-to-device copies with device_copy or thrust::fill_n in the routing cost-boundary t…
Out of Scope Changes check		The PR changes three files. The test change in cost_boundary_initialization.cu serves #2212 directly. The other two changes are production code outside the PDP test path that #2212 names. generator.cu…
Description check		The description explains the change in detail. It says the test copied pickup and delivery indices from stack locals with raft::copy, and that the block ended before any stream sync. The fix builds d_…

Full details: Title check

Explanation

The title names the main change: keeping host sources alive for async copies. The changeset replaces stack-local host-to-device copies with device_copy or thrust::fill_n in the routing cost-boundary test, generator.cu, and problem.cu. The title is specific and matches the changes. It is long at 16 words, but it is not vague or off-topic.

Full details: Out of Scope Changes check

Explanation

The PR changes three files. The test change in cost_boundary_initialization.cu serves #2212 directly. The other two changes are production code outside the PDP test path that #2212 names. generator.cu (generate_coordinates) and problem.cu (csr_to_csc_transpose) replace raft::copy from host locals with thrust::fill_n on the device. The PR states that these sites follow the same hazard pattern as #1759, #2200, and #2202. That connection is explicit and plausible, so the changes are not demonstrably unrelated. The fill_n calls take the scalar by value in the functor, so they do not depend on host source lifetime. The diff is broader than the single test named in #2212, which reviewers may weigh when deciding whether to keep these changes in this PR.

Full details: Description check

Explanation

The description explains the change in detail. It says the test copied pickup and delivery indices from stack locals with raft::copy, and that the block ended before any stream sync. The fix builds d_pickups and d_deliveries with cuopt::device_copy from function-scope vectors. It also covers the depot coordinate change in generator.cu and the CSC offset change in problem.cu, which use thrust::fill_n. The description also lists the testing done and links issue #2212. This matches the three files in the raw summary, so the check passes.

Full details: Docstring Coverage

Explanation

Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff with cyclomatic complexity of at least 2. Documentation-only edits and deletions within surviving functions are included; unchanged diff context and fully deleted functions are excluded. Analyzed 3 functions across 3 files.

Resolution

Write docstrings for the functions missing them to satisfy the coverage threshold.

Full details: Linked Issues check

Explanation

Issue #2212 requires that cost_boundary_initialization.pdp_set_route_data_resets_cost_boundaries pass on gb300. The only coding change for this test is in cpp/tests/routing/unit_tests/cost_boundary_initialization.cu. The PR moves the pickup and delivery host sources (pickups{0}, deliveries{1}) to function scope and builds d_pickups and d_deliveries with cuopt::device_copy. Before, the sources were stack locals inside the if constexpr (is_pdp) block. This removes the source-lifetime hazard in the test. The fix is therefore correct as a lifetime fix. However, the PR does not show that it resolves the gb300 failure. The PR's own probe on B300 and GH200 showed that pageable sources are read at call time. The failure mechanism (driver reads the source after the call returns) was not confirmed on gb300. No gb300 run result is in the evidence. I did not inspect the definition of cuopt::device_copy, so its synchronization behavior is unverified. No unmet coding requirement is demonstrated in the code.

Resolution

Provide the gb300 CI result for cost_boundary_initialization.pdp_set_route_data_resets_cost_boundaries, set_shmem_of_kernel.zero_request, and set_shmem_of_kernel.normal_request on the reviewed head (eba2142). A passing result on gb300 resolves this assessment. Also provide the definition or synchronization semantics of cuopt::device_copy if the reviewer needs to confirm that the function-scope vectors outlive the queued copies.

Finishing Touches

Generate unit tests (beta)

• Create a new PR

• Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

CodeRabbit Verify

• Run CodeRabbit Verify to test the changes in this pull request.

Recent review info

Run configuration

• Source mode: Auto
• Repository: NVIDIA/cuopt
• Branch: main
• File: .coderabbit.yaml
• Commit: 3d71290
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Enterprise
• Run ID: ff9339cc-d4c1-42b9-8ad3-565ec438a7d8

Commits

Reviewing files that changed from the base of the PR and between 3d71290 and eba2142.

Files selected for processing: 3

• cpp/src/mip_heuristics/problem/problem.cu
• cpp/src/routing/generator/generator.cu
• cpp/tests/routing/unit_tests/cost_boundary_initialization.cu

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 9 remain after this review.