Skip to content

Commit

Permalink
fix inappropriate cuco sentinel value when value type is not vertex_t
Browse files Browse the repository at this point in the history
  • Loading branch information
@seunghwak
seunghwak committed Jul 1, 2022
1 parent 0e82187 commit a75dbd0
Show file tree
Hide file tree
Showing 4 changed files with 42 additions and 20 deletions.
17 changes: 11 additions & 6 deletions cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -234,6 +234,11 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e(
VertexIterator map_unique_key_first,
VertexIterator map_unique_key_last,
ValueIterator map_value_first,
#if 1 // FIXME: this is unnecessary if we use a binary tree instead of cuco::static_map in
// collect_values_for_unique_keys, need to compare the two approaches
typename thrust::iterator_traits<VertexIterator>::value_type invalid_key,
typename thrust::iterator_traits<ValueIterator>::value_type invalid_value,
#endif
KeyAggregatedEdgeOp key_aggregated_e_op,
T init,
ReduceOp reduce_op,
Expand Down Expand Up @@ -313,8 +318,8 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e(
static_cast<double>(thrust::distance(map_unique_key_first, map_unique_key_last)) /
load_factor),
static_cast<size_t>(thrust::distance(map_unique_key_first, map_unique_key_last)) + 1),
cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
cuco::sentinel::empty_value<value_t>{0},
cuco::sentinel::empty_key<vertex_t>{invalid_key},
cuco::sentinel::empty_value<value_t>{invalid_value},
stream_adapter,
handle.get_stream());

Expand Down Expand Up @@ -586,8 +591,8 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e(
auto multi_gpu_kv_map_ptr = std::make_unique<
cuco::static_map<vertex_t, value_t, cuda::thread_scope_device, decltype(stream_adapter)>>(
size_t{0},
cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
cuco::sentinel::empty_value<value_t>{0},
cuco::sentinel::empty_key<vertex_t>{invalid_key},
cuco::sentinel::empty_value<value_t>{invalid_value},
stream_adapter,
handle.get_stream()); // relevant only when GraphViewType::is_multi_gpu is true
if constexpr (GraphViewType::is_multi_gpu) {
Expand Down Expand Up @@ -624,8 +629,8 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e(
// cuco::static_map requires at least one empty slot
std::max(static_cast<size_t>(static_cast<double>(unique_minor_keys.size()) / load_factor),
static_cast<size_t>(unique_minor_keys.size()) + 1),
cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
cuco::sentinel::empty_value<value_t>{0},
cuco::sentinel::empty_key<vertex_t>{invalid_key},
cuco::sentinel::empty_value<value_t>{invalid_value},
stream_adapter,
handle.get_stream());

Expand Down
22 changes: 16 additions & 6 deletions cpp/include/cugraph/utilities/collect_comm.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -61,6 +61,11 @@ collect_values_for_keys(raft::comms::comms_t const& comm,
VertexIterator1 collect_key_first,
VertexIterator1 collect_key_last,
KeyToGPUIdOp key_to_gpu_id_op,
#if 1 // FIXME: this is unnecessary if we use a binary tree instead of cuco::static_map, need to
// compare the two approaches
typename thrust::iterator_traits<VertexIterator0>::value_type invalid_key,
typename thrust::iterator_traits<ValueIterator>::value_type invalid_value,
#endif
rmm::cuda_stream_view stream_view)
{
using vertex_t = typename thrust::iterator_traits<VertexIterator0>::value_type;
Expand All @@ -84,8 +89,8 @@ collect_values_for_keys(raft::comms::comms_t const& comm,
std::max(static_cast<size_t>(
static_cast<double>(thrust::distance(map_key_first, map_key_last)) / load_factor),
static_cast<size_t>(thrust::distance(map_key_first, map_key_last)) + 1),
cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
cuco::sentinel::empty_value<value_t>{0},
cuco::sentinel::empty_key<vertex_t>{invalid_key},
cuco::sentinel::empty_value<value_t>{invalid_value},
stream_adapter,
stream_view);
{
Expand Down Expand Up @@ -147,8 +152,8 @@ collect_values_for_keys(raft::comms::comms_t const& comm,
// cuco::static_map requires at least one empty slot
std::max(static_cast<size_t>(static_cast<double>(unique_keys.size()) / load_factor),
unique_keys.size() + 1),
cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
cuco::sentinel::empty_value<value_t>{0},
cuco::sentinel::empty_key<vertex_t>{invalid_key},
cuco::sentinel::empty_value<value_t>{invalid_value},
stream_adapter,
stream_view);
{
Expand Down Expand Up @@ -224,6 +229,11 @@ collect_values_for_unique_keys(
rmm::device_uvector<typename thrust::iterator_traits<VertexIterator>::value_type>&&
collect_unique_keys,
KeyToGPUIdOp key_to_gpu_id_op,
#if 1 // FIXME: this is unnecessary if we use a binary tree instead of cuco::static_map, need to
// compare the two approaches
typename thrust::iterator_traits<VertexIterator>::value_type invalid_key,
typename thrust::iterator_traits<ValueIterator>::value_type invalid_value,
#endif
rmm::cuda_stream_view stream_view)
{
using vertex_t = typename thrust::iterator_traits<VertexIterator>::value_type;
Expand All @@ -244,8 +254,8 @@ collect_values_for_unique_keys(
std::max(static_cast<size_t>(
static_cast<double>(thrust::distance(map_key_first, map_key_last)) / load_factor),
static_cast<size_t>(thrust::distance(map_key_first, map_key_last)) + 1),
cuco::sentinel::empty_key<vertex_t>{invalid_vertex_id<vertex_t>::value},
cuco::sentinel::empty_value<value_t>{0},
cuco::sentinel::empty_key<vertex_t>{invalid_key},
cuco::sentinel::empty_value<value_t>{invalid_value},
stream_adapter,
stream_view);
{
Expand Down
21 changes: 13 additions & 8 deletions cpp/src/community/louvain.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -446,14 +446,17 @@ class Louvain {
cugraph::detail::compute_gpu_id_from_ext_vertex_t<vertex_t> vertex_to_gpu_id_op{
handle_.get_comms().get_size()};

vertex_cluster_weights_v = cugraph::collect_values_for_keys(handle_.get_comms(),
cluster_keys_v_.begin(),
cluster_keys_v_.end(),
cluster_weights_v_.data(),
next_clusters_v_.begin(),
next_clusters_v_.end(),
vertex_to_gpu_id_op,
handle_.get_stream());
vertex_cluster_weights_v =
cugraph::collect_values_for_keys(handle_.get_comms(),
cluster_keys_v_.begin(),
cluster_keys_v_.end(),
cluster_weights_v_.data(),
next_clusters_v_.begin(),
next_clusters_v_.end(),
vertex_to_gpu_id_op,
invalid_vertex_id<vertex_t>::value,
std::numeric_limits<weight_t>::max(),
handle_.get_stream());

src_cluster_weights =
edge_partition_src_property_t<graph_view_t, weight_t>(handle_, current_graph_view_);
Expand Down Expand Up @@ -533,6 +536,8 @@ class Louvain {
cluster_keys_v_.begin(),
cluster_keys_v_.end(),
cluster_weights_v_.begin(),
invalid_vertex_id<vertex_t>::value,
std::numeric_limits<weight_t>::max(),
detail::key_aggregated_edge_op_t<vertex_t, weight_t>{total_edge_weight, resolution},
thrust::make_tuple(vertex_t{-1}, weight_t{0}),
detail::reduce_op_t<vertex_t, weight_t>{},
Expand Down
2 changes: 2 additions & 0 deletions cpp/src/structure/renumber_utils_impl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -410,6 +410,8 @@ void renumber_ext_vertices(raft::handle_t const& handle,
thrust::make_counting_iterator(local_int_vertex_first),
std::move(sorted_unique_ext_vertices),
detail::compute_gpu_id_from_ext_vertex_t<vertex_t>{comm_size},
invalid_vertex_id<vertex_t>::value,
invalid_vertex_id<vertex_t>::value,
handle.get_stream());

renumber_map_ptr.reset();
Expand Down

0 comments on commit a75dbd0

Please sign in to comment.