Skip to content

Commit a7c4ebd

Browse files
authored
Remove literals passed to device_uvector::set_element_async (#1453)
After rapidsai/rmm#725 is merged, this PR updates cuspatial to eliminate passing literal values to device_uvector::set_element_async. Companion PR to rapidsai/cuspatial#367 Authors: - Mark Harris (@harrism) Approvers: - Seunghwa Kang (@seunghwak) - Alex Fender (@afender) - Andrei Schaffer (@aschaffer) URL: #1453
1 parent fe0cfc7 commit a7c4ebd

File tree

1 file changed

+17
-4
lines changed

1 file changed

+17
-4
lines changed

cpp/src/experimental/graph.cu

Lines changed: 17 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -304,9 +304,15 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
304304

305305
rmm::device_uvector<vertex_t> segment_offsets(detail::num_segments_per_vertex_partition + 1,
306306
default_stream);
307-
segment_offsets.set_element_async(0, 0, default_stream);
307+
308+
// temporaries are necessary because the &&-overload of device_uvector is deleted
309+
// Note that we must sync `default_stream` before these temporaries go out of scope to
310+
// avoid use after free. (The syncs are at the end of this function)
311+
auto zero_vertex = vertex_t{0};
312+
auto vertex_count = static_cast<vertex_t>(degrees.size());
313+
segment_offsets.set_element_async(0, zero_vertex, default_stream);
308314
segment_offsets.set_element_async(
309-
detail::num_segments_per_vertex_partition, degrees.size(), default_stream);
315+
detail::num_segments_per_vertex_partition, vertex_count, default_stream);
310316

311317
thrust::upper_bound(rmm::exec_policy(default_stream)->on(default_stream),
312318
degrees.begin(),
@@ -454,9 +460,16 @@ graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enable_if_
454460

455461
rmm::device_uvector<vertex_t> segment_offsets(detail::num_segments_per_vertex_partition + 1,
456462
default_stream);
457-
segment_offsets.set_element_async(0, 0, default_stream);
463+
464+
// temporaries are necessary because the &&-overload of device_uvector is deleted
465+
// Note that we must sync `default_stream` before these temporaries go out of scope to
466+
// avoid use after free. (The syncs are at the end of this function)
467+
auto zero_vertex = vertex_t{0};
468+
auto vertex_count = static_cast<vertex_t>(this->get_number_of_vertices());
469+
segment_offsets.set_element_async(0, zero_vertex, default_stream);
470+
458471
segment_offsets.set_element_async(
459-
detail::num_segments_per_vertex_partition, this->get_number_of_vertices(), default_stream);
472+
detail::num_segments_per_vertex_partition, vertex_count, default_stream);
460473

461474
thrust::upper_bound(rmm::exec_policy(default_stream)->on(default_stream),
462475
degree_first,

0 commit comments

Comments
 (0)