Skip to content

Keep host sources alive for async copies in the routing cost-boundary test and two src sites #2260

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
0z5a wants to merge 2 commits into
base: main
Choose a base branch
from
+12 −9
Open

Keep host sources alive for async copies in the routing cost-boundary test and two src sites #2260

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
da17c19
Keep host sources alive for async copies in the cost-boundary test an…
0z5a Nov 3, 2026
eba2142
Copy the PDP pickup/delivery indices only in the PDP variant
0z5a Nov 3, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion cpp/src/mip_heuristics/problem/problem.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@
#include <thrust/binary_search.h>
#include <thrust/copy.h>
#include <thrust/count.h>
#include <thrust/fill.h>
#include <thrust/gather.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
Expand Down Expand Up @@ -466,7 +467,7 @@ void csr_to_csc_transpose(const i_t* csr_offsets,
raft::copy(&last_ptr, csc_offsets + (n_cols - 1), 1, stream);
handle_ptr->sync_stream();
i_t total_nnz = last_ptr + last_count;
raft::copy(csc_offsets + n_cols, &total_nnz, 1, stream);
thrust::fill_n(handle_ptr->get_thrust_policy(), csc_offsets + n_cols, 1, total_nnz);

@ramakrishnap-nv ramakrishnap-nv Nov 3, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 0z5a Nov 3, 2026

Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.


// Scatter values
rmm::device_uvector<i_t> next_pos(n_cols, stream);
Expand Down
5 changes: 3 additions & 2 deletions cpp/src/routing/generator/generator.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@
#include "generator.hpp"
#include "generator_utils.cuh"

#include <thrust/fill.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/reduce.h>
#include <thrust/shuffle.h>
Expand Down Expand Up @@ -207,8 +208,8 @@ coordinates_t<f_t> generate_coordinates(raft::handle_t& handle,
auto depot_y = depot_x;

// Update depot to be in the center.
raft::copy(v_x_pos.data(), &depot_x, 1, handle.get_stream());
raft::copy(v_y_pos.data(), &depot_y, 1, handle.get_stream());
thrust::fill_n(handle.get_thrust_policy(), v_x_pos.begin(), 1, depot_x);

@ramakrishnap-nv ramakrishnap-nv Nov 3, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 0z5a Nov 3, 2026 •
edited
Loading

Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

thrust::fill_n(handle.get_thrust_policy(), v_y_pos.begin(), 1, depot_y);
return std::make_tuple(std::move(v_x_pos), std::move(v_y_pos));
}

Expand Down
13 changes: 7 additions & 6 deletions cpp/tests/routing/unit_tests/cost_boundary_initialization.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,13 +59,14 @@ void test_cost_boundaries_reset()
data_model.add_cost_matrix(d_cost_matrix.data());
data_model.set_order_locations(d_order_locations.data());

rmm::device_uvector<int> d_pickups(is_pdp ? 1 : 0, stream);
rmm::device_uvector<int> d_deliveries(is_pdp ? 1 : 0, stream);
// Function scope, so the host sources outlive the asynchronous copies below.
const std::vector<int> pickups{0};
const std::vector<int> deliveries{1};
rmm::device_uvector<int> d_pickups(0, stream);
rmm::device_uvector<int> d_deliveries(0, stream);
if constexpr (is_pdp) {
const int pickup = 0;
const int delivery = 1;
raft::copy(d_pickups.data(), &pickup, 1, stream);
raft::copy(d_deliveries.data(), &delivery, 1, stream);
d_pickups = cuopt::device_copy(pickups, stream);
d_deliveries = cuopt::device_copy(deliveries, stream);
data_model.set_pickup_delivery_pairs(d_pickups.data(), d_deliveries.data());
}

Expand Down