Skip to content

Commit

Permalink
Separate edge weights from graph objects and update primitives to sup…
Browse files Browse the repository at this point in the history 
…port general edge properties. (#2843)

This PR separates edge weights from `graph_t` and `graph_view_t`. Now edge weights are stored in `edge_property_t` like any other edge properties (e.g. edge IDs & types).

Two major benefits are

1) For algorithms working on unweighted graphs, we don't need to explicitly instantiate for weight_t = float and weight_t = double cases. This saves compile time & reduces the binary size.
2) Primitives can work with different edge properties (weights, IDs, types, and so on) using a single mechanism instead of using one mechanism for edge weights and another mechanism for other edge properties (e.g. edge IDs & types).

This PR changes many files (due to the changes in the `graph_t` and `graph_view_t` classes) but most changes are repetitive.

To highlight major changes,

1)
```
template <typename vertex_t, typename edge_t, typename weight_t, bool store_transposed, bool multi_gpu>
class graph_(view_)t {
...
}
=>
template <typename vertex_t, typename edge_t, bool store_transposed, bool multi_gpu>
class graph_(view_)t {
...
}
```

2)
`create_graph_from_edgelist` will return 
```
std::tuple<
  graph_t<vertex_t, edge_t, store_transposed, multi_gpu>,
  std::optional<
    edge_property_t<graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu>, weight_t>>,
  std::optional<edge_property_t<graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu>,
                                thrust::tuple<edge_t, edge_type_t>>>,
  std::optional<rmm::device_uvector<vertex_t>>>
```
instead of
```
std::tuple<
  graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu>,
  std::optional<edge_property_t<graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu>,
                                thrust::tuple<edge_t, edge_type_t>>>,
  std::optional<rmm::device_uvector<vertex_t>>>
```
i.e. edge weights will be returned in a separate object.

3)
Algorithms working on weighted graphs (or both unweighted & weighted graphs) will take the additional input argument:
`edge_property_view_t<edge_t, weight_t const*>` or `std::optional<edge_property_view_t<edge_t, weight_t const*>>`

4)
Primitives `e_op`'s input parameters were previously `(vertex_t src, vertex_t dst, weight_t w, source_property, destination_property)` or `(vertex_t src, vertex_t dst, source_property, destination_property)`.
Now it will be `(vertex_t src, vertex_t dst, source_property, destination property, edge property)`
Primitives with `e_op` will take an additional `edge_property_view_t` type input parameter (which can be edge_dummy_property_t{}.view() if no edge properties are used).

5)
Edge weights are added in the C++ side (before type erasing) of the C API `cugraph_graph_t`. 
```
struct cugraph_graph_t {
...
  void* edge_weights_;  // this is added.
...
}
```

This doesn't require any changes in the C side of the C API, so I assume this won't disrupt the pylibcugraph and python layers (once cython.cu file is removed).

Authors:
  - Seunghwa Kang (https://github.com/seunghwak)

Approvers:
  - Chuck Hastings (https://github.com/ChuckHastings)
  - Naim (https://github.com/naimnv)

URL: #2843
  • Loading branch information
@seunghwak
seunghwak authored Nov 21, 2022
1 parent aab7103 commit 4b7d4fb
Show file tree
Hide file tree
Showing 241 changed files with 6,890 additions and 8,230 deletions.
273 changes: 154 additions & 119 deletions cpp/include/cugraph/algorithms.hpp
View file

Large diffs are not rendered by default.

358 changes: 19 additions & 339 deletions cpp/include/cugraph/detail/decompress_edge_partition.cuh
View file

Large diffs are not rendered by default.

57 changes: 16 additions & 41 deletions cpp/include/cugraph/edge_partition_device_view.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -92,16 +92,12 @@ struct local_degree_op_t {
}
};

template <typename vertex_t, typename edge_t, typename weight_t>
template <typename vertex_t, typename edge_t>
class edge_partition_device_view_base_t {
public:
edge_partition_device_view_base_t(raft::device_span<edge_t const> offsets,
raft::device_span<vertex_t const> indices,
std::optional<raft::device_span<weight_t const>> weights)
: offsets_(offsets),
indices_(indices),
weights_(weights ? thrust::optional<raft::device_span<weight_t const>>(*weights)
: thrust::nullopt)
raft::device_span<vertex_t const> indices)
: offsets_(offsets), indices_(indices)
{
}

Expand All @@ -112,21 +108,15 @@ class edge_partition_device_view_base_t {

__host__ __device__ edge_t const* offsets() const { return offsets_.data(); }
__host__ __device__ vertex_t const* indices() const { return indices_.data(); }
__host__ __device__ thrust::optional<weight_t const*> weights() const
{
return weights_ ? thrust::optional<weight_t const*>{(*weights_).data()} : thrust::nullopt;
}

// major_idx == major offset if CSR/CSC, major_offset != major_idx if DCSR/DCSC
__device__ thrust::tuple<vertex_t const*, thrust::optional<weight_t const*>, edge_t> local_edges(
__device__ thrust::tuple<vertex_t const*, edge_t, edge_t> local_edges(
vertex_t major_idx) const noexcept
{
auto edge_offset = offsets_[major_idx];
auto local_degree = offsets_[major_idx + 1] - edge_offset;
auto indices = indices_.data() + edge_offset;
auto weights = weights_ ? thrust::optional<weight_t const*>{(*weights_).data() + edge_offset}
: thrust::nullopt;
return thrust::make_tuple(indices, weights, local_degree);
return thrust::make_tuple(indices, edge_offset, local_degree);
}

// major_idx == major offset if CSR/CSC, major_offset != major_idx if DCSR/DCSC
Expand All @@ -142,30 +132,20 @@ class edge_partition_device_view_base_t {
// should be trivially copyable to device
raft::device_span<edge_t const> offsets_{nullptr};
raft::device_span<vertex_t const> indices_{nullptr};
thrust::optional<raft::device_span<weight_t const>> weights_{thrust::nullopt};
};

} // namespace detail

template <typename vertex_t,
typename edge_t,
typename weight_t,
bool multi_gpu,
typename Enable = void>
template <typename vertex_t, typename edge_t, bool multi_gpu, typename Enable = void>
class edge_partition_device_view_t;

// multi-GPU version
template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
class edge_partition_device_view_t<vertex_t,
edge_t,
weight_t,
multi_gpu,
std::enable_if_t<multi_gpu>>
: public detail::edge_partition_device_view_base_t<vertex_t, edge_t, weight_t> {
template <typename vertex_t, typename edge_t, bool multi_gpu>
class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t<multi_gpu>>
: public detail::edge_partition_device_view_base_t<vertex_t, edge_t> {
public:
edge_partition_device_view_t(edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu> view)
: detail::edge_partition_device_view_base_t<vertex_t, edge_t, weight_t>(
view.offsets(), view.indices(), view.weights()),
edge_partition_device_view_t(edge_partition_view_t<vertex_t, edge_t, multi_gpu> view)
: detail::edge_partition_device_view_base_t<vertex_t, edge_t>(view.offsets(), view.indices()),
dcs_nzd_vertices_(detail::to_thrust_optional(view.dcs_nzd_vertices())),
major_hypersparse_first_(detail::to_thrust_optional(view.major_hypersparse_first())),
major_range_first_(view.major_range_first()),
Expand Down Expand Up @@ -359,17 +339,12 @@ class edge_partition_device_view_t<vertex_t,
};

// single-GPU version
template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
class edge_partition_device_view_t<vertex_t,
edge_t,
weight_t,
multi_gpu,
std::enable_if_t<!multi_gpu>>
: public detail::edge_partition_device_view_base_t<vertex_t, edge_t, weight_t> {
template <typename vertex_t, typename edge_t, bool multi_gpu>
class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t<!multi_gpu>>
: public detail::edge_partition_device_view_base_t<vertex_t, edge_t> {
public:
edge_partition_device_view_t(edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu> view)
: detail::edge_partition_device_view_base_t<vertex_t, edge_t, weight_t>(
view.offsets(), view.indices(), view.weights()),
edge_partition_device_view_t(edge_partition_view_t<vertex_t, edge_t, multi_gpu> view)
: detail::edge_partition_device_view_base_t<vertex_t, edge_t>(view.offsets(), view.indices()),
number_of_vertices_(view.major_range_last())
{
}
Expand Down
2 changes: 2 additions & 0 deletions cpp/include/cugraph/edge_partition_edge_property_device_view.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,8 @@ class edge_partition_edge_property_device_view_t {
value_first_ = view.value_firsts()[partition_idx];
}

__host__ __device__ ValueIterator value_first() { return value_first_; }

__device__ ValueIterator get_iter(edge_t offset) const { return value_first_ + offset; }

__device__ value_type get(edge_t offset) const { return *get_iter(offset); }
Expand Down
33 changes: 12 additions & 21 deletions cpp/include/cugraph/edge_partition_view.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,53 +24,45 @@ namespace cugraph {

namespace detail {

template <typename vertex_t, typename edge_t, typename weight_t>
template <typename vertex_t, typename edge_t>
class edge_partition_view_base_t {
public:
edge_partition_view_base_t(raft::device_span<edge_t const> offsets,
raft::device_span<vertex_t const> indices,
std::optional<raft::device_span<weight_t const>> weights)
: offsets_(offsets), indices_(indices), weights_(weights)
raft::device_span<vertex_t const> indices)
: offsets_(offsets), indices_(indices)
{
}

edge_t number_of_edges() const { return static_cast<edge_t>(indices_.size()); }

raft::device_span<edge_t const> offsets() const { return offsets_; }
raft::device_span<vertex_t const> indices() const { return indices_; }
std::optional<raft::device_span<weight_t const>> weights() const { return weights_; }

private:
raft::device_span<edge_t const> offsets_{};
raft::device_span<vertex_t const> indices_{};
std::optional<raft::device_span<weight_t const>> weights_{std::nullopt};
};

} // namespace detail

template <typename vertex_t,
typename edge_t,
typename weight_t,
bool multi_gpu,
typename Enable = void>
template <typename vertex_t, typename edge_t, bool multi_gpu, typename Enable = void>
class edge_partition_view_t;

// multi-GPU version
template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_if_t<multi_gpu>>
: public detail::edge_partition_view_base_t<vertex_t, edge_t, weight_t> {
template <typename vertex_t, typename edge_t, bool multi_gpu>
class edge_partition_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t<multi_gpu>>
: public detail::edge_partition_view_base_t<vertex_t, edge_t> {
public:
edge_partition_view_t(raft::device_span<edge_t const> offsets,
raft::device_span<vertex_t const> indices,
std::optional<raft::device_span<weight_t const>> weights,
std::optional<raft::device_span<vertex_t const>> dcs_nzd_vertices,
std::optional<vertex_t> major_hypersparse_first,
vertex_t major_range_first,
vertex_t major_range_last,
vertex_t minor_range_first,
vertex_t minor_range_last,
vertex_t major_value_start_offset)
: detail::edge_partition_view_base_t<vertex_t, edge_t, weight_t>(offsets, indices, weights),
: detail::edge_partition_view_base_t<vertex_t, edge_t>(offsets, indices),
dcs_nzd_vertices_(dcs_nzd_vertices),
major_hypersparse_first_(major_hypersparse_first),
major_range_first_(major_range_first),
Expand Down Expand Up @@ -109,15 +101,14 @@ class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_i
};

// single-GPU version
template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
class edge_partition_view_t<vertex_t, edge_t, weight_t, multi_gpu, std::enable_if_t<!multi_gpu>>
: public detail::edge_partition_view_base_t<vertex_t, edge_t, weight_t> {
template <typename vertex_t, typename edge_t, bool multi_gpu>
class edge_partition_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t<!multi_gpu>>
: public detail::edge_partition_view_base_t<vertex_t, edge_t> {
public:
edge_partition_view_t(raft::device_span<edge_t const> offsets,
raft::device_span<vertex_t const> indices,
std::optional<raft::device_span<weight_t const>> weights,
vertex_t number_of_vertices)
: detail::edge_partition_view_base_t<vertex_t, edge_t, weight_t>(offsets, indices, weights),
: detail::edge_partition_view_base_t<vertex_t, edge_t>(offsets, indices),
number_of_vertices_(number_of_vertices)
{
}
Expand Down
14 changes: 5 additions & 9 deletions cpp/include/cugraph/edge_property.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -27,8 +27,6 @@

namespace cugraph {

namespace detail {

template <typename edge_t, typename ValueIterator>
class edge_property_view_t {
public:
Expand Down Expand Up @@ -59,8 +57,6 @@ class edge_dummy_property_view_t {
using value_iterator = void*;
};

} // namespace detail

template <typename GraphViewType, typename T>
class edge_property_t {
public:
Expand Down Expand Up @@ -98,8 +94,8 @@ class edge_property_t {
edge_partition_edge_counts[i] = size_dataframe_buffer(buffers_[i]);
}

return detail::edge_property_view_t<edge_type, const_value_iterator>(
edge_partition_value_firsts, edge_partition_edge_counts);
return edge_property_view_t<edge_type, const_value_iterator>(edge_partition_value_firsts,
edge_partition_edge_counts);
}

auto mutable_view()
Expand All @@ -113,8 +109,8 @@ class edge_property_t {
edge_partition_edge_counts[i] = size_dataframe_buffer(buffers_[i]);
}

return detail::edge_property_view_t<edge_type, value_iterator>(edge_partition_value_firsts,
edge_partition_edge_counts);
return edge_property_view_t<edge_type, value_iterator>(edge_partition_value_firsts,
edge_partition_edge_counts);
}

private:
Expand All @@ -125,7 +121,7 @@ class edge_dummy_property_t {
public:
using value_type = thrust::nullopt_t;

auto view() const { return detail::edge_dummy_property_view_t{}; }
auto view() const { return edge_dummy_property_view_t{}; }
};

} // namespace cugraph
98 changes: 33 additions & 65 deletions cpp/include/cugraph/eidecl_graph.hpp
View file
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020-2021, NVIDIA CORPORATION.
* Copyright (c) 2020-2022, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand All @@ -16,68 +16,36 @@
#pragma once

namespace cugraph {
extern template class graph_t<int32_t, int32_t, float, true, true, void>;
extern template class graph_t<int32_t, int32_t, float, true, false, void>;
extern template class graph_t<int32_t, int32_t, float, false, true, void>;
extern template class graph_t<int32_t, int32_t, float, false, false, void>;
extern template class graph_t<int32_t, int32_t, double, true, true, void>;
extern template class graph_t<int32_t, int32_t, double, true, false, void>;
extern template class graph_t<int32_t, int32_t, double, false, true, void>;
extern template class graph_t<int32_t, int32_t, double, false, false, void>;
extern template class graph_t<int32_t, int64_t, float, true, true, void>;
extern template class graph_t<int32_t, int64_t, float, true, false, void>;
extern template class graph_t<int32_t, int64_t, float, false, true, void>;
extern template class graph_t<int32_t, int64_t, float, false, false, void>;
extern template class graph_t<int32_t, int64_t, double, true, true, void>;
extern template class graph_t<int32_t, int64_t, double, true, false, void>;
extern template class graph_t<int32_t, int64_t, double, false, true, void>;
extern template class graph_t<int32_t, int64_t, double, false, false, void>;
extern template class graph_t<int64_t, int32_t, float, true, true, void>;
extern template class graph_t<int64_t, int32_t, float, true, false, void>;
extern template class graph_t<int64_t, int32_t, float, false, true, void>;
extern template class graph_t<int64_t, int32_t, float, false, false, void>;
extern template class graph_t<int64_t, int32_t, double, true, true, void>;
extern template class graph_t<int64_t, int32_t, double, true, false, void>;
extern template class graph_t<int64_t, int32_t, double, false, true, void>;
extern template class graph_t<int64_t, int32_t, double, false, false, void>;
extern template class graph_t<int64_t, int64_t, float, true, true, void>;
extern template class graph_t<int64_t, int64_t, float, true, false, void>;
extern template class graph_t<int64_t, int64_t, float, false, true, void>;
extern template class graph_t<int64_t, int64_t, float, false, false, void>;
extern template class graph_t<int64_t, int64_t, double, true, true, void>;
extern template class graph_t<int64_t, int64_t, double, true, false, void>;
extern template class graph_t<int64_t, int64_t, double, false, true, void>;
extern template class graph_t<int64_t, int64_t, double, false, false, void>;
extern template class graph_view_t<int32_t, int32_t, float, true, true, void>;
extern template class graph_view_t<int32_t, int32_t, float, true, false, void>;
extern template class graph_view_t<int32_t, int32_t, float, false, true, void>;
extern template class graph_view_t<int32_t, int32_t, float, false, false, void>;
extern template class graph_view_t<int32_t, int32_t, double, true, true, void>;
extern template class graph_view_t<int32_t, int32_t, double, true, false, void>;
extern template class graph_view_t<int32_t, int32_t, double, false, true, void>;
extern template class graph_view_t<int32_t, int32_t, double, false, false, void>;
extern template class graph_view_t<int32_t, int64_t, float, true, true, void>;
extern template class graph_view_t<int32_t, int64_t, float, true, false, void>;
extern template class graph_view_t<int32_t, int64_t, float, false, true, void>;
extern template class graph_view_t<int32_t, int64_t, float, false, false, void>;
extern template class graph_view_t<int32_t, int64_t, double, true, true, void>;
extern template class graph_view_t<int32_t, int64_t, double, true, false, void>;
extern template class graph_view_t<int32_t, int64_t, double, false, true, void>;
extern template class graph_view_t<int32_t, int64_t, double, false, false, void>;
extern template class graph_view_t<int64_t, int32_t, float, true, true, void>;
extern template class graph_view_t<int64_t, int32_t, float, true, false, void>;
extern template class graph_view_t<int64_t, int32_t, float, false, true, void>;
extern template class graph_view_t<int64_t, int32_t, float, false, false, void>;
extern template class graph_view_t<int64_t, int32_t, double, true, true, void>;
extern template class graph_view_t<int64_t, int32_t, double, true, false, void>;
extern template class graph_view_t<int64_t, int32_t, double, false, true, void>;
extern template class graph_view_t<int64_t, int32_t, double, false, false, void>;
extern template class graph_view_t<int64_t, int64_t, float, true, true, void>;
extern template class graph_view_t<int64_t, int64_t, float, true, false, void>;
extern template class graph_view_t<int64_t, int64_t, float, false, true, void>;
extern template class graph_view_t<int64_t, int64_t, float, false, false, void>;
extern template class graph_view_t<int64_t, int64_t, double, true, true, void>;
extern template class graph_view_t<int64_t, int64_t, double, true, false, void>;
extern template class graph_view_t<int64_t, int64_t, double, false, true, void>;
extern template class graph_view_t<int64_t, int64_t, double, false, false, void>;
extern template class graph_t<int32_t, int32_t, true, true, void>;
extern template class graph_t<int32_t, int32_t, true, false, void>;
extern template class graph_t<int32_t, int32_t, false, true, void>;
extern template class graph_t<int32_t, int32_t, false, false, void>;
extern template class graph_t<int32_t, int64_t, true, true, void>;
extern template class graph_t<int32_t, int64_t, true, false, void>;
extern template class graph_t<int32_t, int64_t, false, true, void>;
extern template class graph_t<int32_t, int64_t, false, false, void>;
extern template class graph_t<int64_t, int32_t, true, true, void>;
extern template class graph_t<int64_t, int32_t, true, false, void>;
extern template class graph_t<int64_t, int32_t, false, true, void>;
extern template class graph_t<int64_t, int32_t, false, false, void>;
extern template class graph_t<int64_t, int64_t, true, true, void>;
extern template class graph_t<int64_t, int64_t, true, false, void>;
extern template class graph_t<int64_t, int64_t, false, true, void>;
extern template class graph_t<int64_t, int64_t, false, false, void>;
extern template class graph_view_t<int32_t, int32_t, true, true, void>;
extern template class graph_view_t<int32_t, int32_t, true, false, void>;
extern template class graph_view_t<int32_t, int32_t, false, true, void>;
extern template class graph_view_t<int32_t, int32_t, false, false, void>;
extern template class graph_view_t<int32_t, int64_t, true, true, void>;
extern template class graph_view_t<int32_t, int64_t, true, false, void>;
extern template class graph_view_t<int32_t, int64_t, false, true, void>;
extern template class graph_view_t<int32_t, int64_t, false, false, void>;
extern template class graph_view_t<int64_t, int32_t, true, true, void>;
extern template class graph_view_t<int64_t, int32_t, true, false, void>;
extern template class graph_view_t<int64_t, int32_t, false, true, void>;
extern template class graph_view_t<int64_t, int32_t, false, false, void>;
extern template class graph_view_t<int64_t, int64_t, true, true, void>;
extern template class graph_view_t<int64_t, int64_t, true, false, void>;
extern template class graph_view_t<int64_t, int64_t, false, true, void>;
extern template class graph_view_t<int64_t, int64_t, false, false, void>;
} // namespace cugraph
Loading

0 comments on commit 4b7d4fb

Please sign in to comment.