Merge branch 'branch-24.02' into drop_support_for_pascal

ci/build_wheel_cugraph.sh

@@ -12,6 +12,6 @@ RAPIDS_PY_CUDA_SUFFIX="$(rapids-wheel-ctk-name-gen ${RAPIDS_CUDA_VERSION})"
 RAPIDS_PY_WHEEL_NAME=pylibcugraph_${RAPIDS_PY_CUDA_SUFFIX} rapids-download-wheels-from-s3 ./local-pylibcugraph
 export PIP_FIND_LINKS=$(pwd)/local-pylibcugraph
 
-export SKBUILD_CONFIGURE_OPTIONS="-DDETECT_CONDA_ENV=OFF -DCUGRAPH_BUILD_WHEELS=ON -DFIND_CUGRAPH_CPP=OFF -DCPM_cugraph-ops_SOURCE=${GITHUB_WORKSPACE}/cugraph-ops/"
+export SKBUILD_CONFIGURE_OPTIONS="-DDETECT_CONDA_ENV=OFF -DFIND_CUGRAPH_CPP=OFF -DCPM_cugraph-ops_SOURCE=${GITHUB_WORKSPACE}/cugraph-ops/"
 
 ./ci/build_wheel.sh cugraph python/cugraph

ci/build_wheel_pylibcugraph.sh

@@ -3,6 +3,6 @@
 
 set -euo pipefail
 
-export SKBUILD_CONFIGURE_OPTIONS="-DDETECT_CONDA_ENV=OFF -DCUGRAPH_BUILD_WHEELS=ON -DFIND_CUGRAPH_CPP=OFF -DCPM_cugraph-ops_SOURCE=${GITHUB_WORKSPACE}/cugraph-ops/"
+export SKBUILD_CONFIGURE_OPTIONS="-DDETECT_CONDA_ENV=OFF -DFIND_CUGRAPH_CPP=OFF -DCPM_cugraph-ops_SOURCE=${GITHUB_WORKSPACE}/cugraph-ops/"
 
 ./ci/build_wheel.sh pylibcugraph python/pylibcugraph

cpp/CMakeLists.txt

@@ -423,7 +423,7 @@ add_library(cugraph_c
         src/c_api/core_result.cpp
         src/c_api/extract_ego.cpp
         src/c_api/k_core.cpp
-	src/c_api/hierarchical_clustering_result.cpp
+        src/c_api/hierarchical_clustering_result.cpp
         src/c_api/induced_subgraph.cpp
         src/c_api/capi_helper.cu
         src/c_api/legacy_spectral.cpp

cpp/include/cugraph/graph.hpp

@@ -90,24 +90,25 @@ class graph_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if_t<mu
     graph_meta_t<vertex_t, edge_t, multi_gpu> meta,
     bool do_expensive_check = false);
 
+  edge_t number_of_edges() const { return this->number_of_edges_; }
+
   graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu> view() const
   {
-    std::vector<edge_t const*> offsets(edge_partition_offsets_.size(), nullptr);
-    std::vector<vertex_t const*> indices(edge_partition_indices_.size(), nullptr);
-    auto dcs_nzd_vertices      = edge_partition_dcs_nzd_vertices_
-                                   ? std::make_optional<std::vector<vertex_t const*>>(
-                                  (*edge_partition_dcs_nzd_vertices_).size(), nullptr)
-                                   : std::nullopt;
-    auto dcs_nzd_vertex_counts = edge_partition_dcs_nzd_vertex_counts_
-                                   ? std::make_optional<std::vector<vertex_t>>(
-                                       (*edge_partition_dcs_nzd_vertex_counts_).size(), vertex_t{0})
-                                   : std::nullopt;
+    std::vector<raft::device_span<edge_t const>> offsets(edge_partition_offsets_.size());
+    std::vector<raft::device_span<vertex_t const>> indices(edge_partition_indices_.size());
+    auto dcs_nzd_vertices = edge_partition_dcs_nzd_vertices_
+                              ? std::make_optional<std::vector<raft::device_span<vertex_t const>>>(
+                                  (*edge_partition_dcs_nzd_vertices_).size())
+                              : std::nullopt;
     for (size_t i = 0; i < offsets.size(); ++i) {
-      offsets[i] = edge_partition_offsets_[i].data();
-      indices[i] = edge_partition_indices_[i].data();
+      offsets[i] = raft::device_span<edge_t const>(edge_partition_offsets_[i].data(),
+                                                   edge_partition_offsets_[i].size());
+      indices[i] = raft::device_span<vertex_t const>(edge_partition_indices_[i].data(),
+                                                     edge_partition_indices_[i].size());
       if (dcs_nzd_vertices) {
-        (*dcs_nzd_vertices)[i]      = (*edge_partition_dcs_nzd_vertices_)[i].data();
-        (*dcs_nzd_vertex_counts)[i] = (*edge_partition_dcs_nzd_vertex_counts_)[i];
+        (*dcs_nzd_vertices)[i] =
+          raft::device_span<vertex_t const>((*edge_partition_dcs_nzd_vertices_)[i].data(),
+                                            (*edge_partition_dcs_nzd_vertices_)[i].size());
       }
     }
 
@@ -196,15 +197,13 @@ class graph_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if_t<mu
     }
 
     return graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu>(
-      *(this->handle_ptr()),
       offsets,
       indices,
       dcs_nzd_vertices,
-      dcs_nzd_vertex_counts,
       graph_view_meta_t<vertex_t, edge_t, store_transposed, multi_gpu>{
         this->number_of_vertices(),
         this->number_of_edges(),
-        this->graph_properties(),
+        this->properties_,
         partition_,
         edge_partition_segment_offsets_,
         local_sorted_unique_edge_srcs,
@@ -224,7 +223,6 @@ class graph_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if_t<mu
   // nzd: nonzero (local) degree
   std::optional<std::vector<rmm::device_uvector<vertex_t>>> edge_partition_dcs_nzd_vertices_{
     std::nullopt};
-  std::optional<std::vector<vertex_t>> edge_partition_dcs_nzd_vertex_counts_{std::nullopt};
   partition_t<vertex_t> partition_{};
 
   // segment offsets within the vertex partition based on vertex degree
@@ -283,16 +281,15 @@ class graph_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if_t<!m
           graph_meta_t<vertex_t, edge_t, multi_gpu> meta,
           bool do_expensive_check = false);
 
+  edge_t number_of_edges() const { return this->number_of_edges_; }
+
   graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu> view() const
   {
     return graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu>(
-      *(this->handle_ptr()),
-      offsets_.data(),
-      indices_.data(),
-      graph_view_meta_t<vertex_t, edge_t, store_transposed, multi_gpu>{this->number_of_vertices(),
-                                                                       this->number_of_edges(),
-                                                                       this->graph_properties(),
-                                                                       segment_offsets_});
+      raft::device_span<edge_t const>(offsets_.data(), offsets_.size()),
+      raft::device_span<vertex_t const>(indices_.data(), indices_.size()),
+      graph_view_meta_t<vertex_t, edge_t, store_transposed, multi_gpu>{
+        this->number_of_vertices(), this->number_of_edges(), this->properties_, segment_offsets_});
   }
 
  private:

cpp/include/cugraph/graph_view.hpp

@@ -258,21 +258,12 @@ class graph_base_t {
  public:
   graph_base_t() = default;
 
-  graph_base_t(raft::handle_t const& handle,
-               vertex_t number_of_vertices,
-               edge_t number_of_edges,
-               graph_properties_t properties)
-    : handle_ptr_(&handle),
-      number_of_vertices_(number_of_vertices),
+  graph_base_t(vertex_t number_of_vertices, edge_t number_of_edges, graph_properties_t properties)
+    : number_of_vertices_(number_of_vertices),
       number_of_edges_(number_of_edges),
       properties_(properties){};
 
   vertex_t number_of_vertices() const { return number_of_vertices_; }
-  edge_t number_of_edges() const
-  {
-    CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented.");
-    return number_of_edges_;
-  }
 
   template <typename vertex_type = vertex_t>
   std::enable_if_t<std::is_signed<vertex_type>::value, bool> is_valid_vertex(vertex_type v) const
@@ -289,33 +280,12 @@ class graph_base_t {
   bool is_symmetric() const { return properties_.is_symmetric; }
   bool is_multigraph() const { return properties_.is_multigraph; }
 
-  void attach_edge_mask(edge_property_view_t<edge_t, uint32_t const*, bool> edge_mask_view)
-  {
-    edge_mask_view_ = edge_mask_view;
-  }
-
-  void clear_edge_mask() { edge_mask_view_ = std::nullopt; }
-
-  bool has_edge_mask() const { return edge_mask_view_.has_value(); }
-
-  std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view() const
-  {
-    return edge_mask_view_;
-  }
-
  protected:
-  raft::handle_t const* handle_ptr() const { return handle_ptr_; };
-  graph_properties_t graph_properties() const { return properties_; }
-
- private:
-  raft::handle_t const* handle_ptr_{nullptr};
-
-  vertex_t number_of_vertices_{0};
   edge_t number_of_edges_{0};
-
   graph_properties_t properties_{};
 
-  std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view_{std::nullopt};
+ private:
+  vertex_t number_of_vertices_{0};
 };
 
 }  // namespace detail
@@ -405,11 +375,10 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
   static constexpr bool is_storage_transposed = store_transposed;
   static constexpr bool is_multi_gpu          = multi_gpu;
 
-  graph_view_t(raft::handle_t const& handle,
-               std::vector<edge_t const*> const& edge_partition_offsets,
-               std::vector<vertex_t const*> const& edge_partition_indices,
-               std::optional<std::vector<vertex_t const*>> const& edge_partition_dcs_nzd_vertices,
-               std::optional<std::vector<vertex_t>> const& edge_partition_dcs_nzd_vertex_counts,
+  graph_view_t(std::vector<raft::device_span<edge_t const>> const& edge_partition_offsets,
+               std::vector<raft::device_span<vertex_t const>> const& edge_partition_indices,
+               std::optional<std::vector<raft::device_span<vertex_t const>>> const&
+                 edge_partition_dcs_nzd_vertices,
                graph_view_meta_t<vertex_t, edge_t, store_transposed, multi_gpu> meta);
 
   std::vector<vertex_t> vertex_partition_range_offsets() const
@@ -624,25 +593,16 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
         this->local_edge_partition_src_value_start_offset(partition_idx);
     }
     std::optional<vertex_t> major_hypersparse_first{std::nullopt};
-    vertex_t offset_size = (major_range_last - major_range_first) + 1;
     if (this->use_dcs()) {
       major_hypersparse_first =
         major_range_first + (*(this->local_edge_partition_segment_offsets(
                               partition_idx)))[detail::num_sparse_segments_per_vertex_partition];
-      offset_size = ((*major_hypersparse_first) - major_range_first) +
-                    (*edge_partition_dcs_nzd_vertex_counts_)[partition_idx] + 1;
     }
     return edge_partition_view_t<vertex_t, edge_t, true>(
-      raft::device_span<edge_t const>(edge_partition_offsets_[partition_idx],
-                                      edge_partition_offsets_[partition_idx] + offset_size),
-      raft::device_span<vertex_t const>(
-        edge_partition_indices_[partition_idx],
-        edge_partition_indices_[partition_idx] + edge_partition_number_of_edges_[partition_idx]),
+      edge_partition_offsets_[partition_idx],
+      edge_partition_indices_[partition_idx],
       edge_partition_dcs_nzd_vertices_
-        ? std::make_optional<raft::device_span<vertex_t const>>(
-            (*edge_partition_dcs_nzd_vertices_)[partition_idx],
-            (*edge_partition_dcs_nzd_vertices_)[partition_idx] +
-              (*edge_partition_dcs_nzd_vertex_counts_)[partition_idx])
+        ? std::make_optional((*edge_partition_dcs_nzd_vertices_)[partition_idx])
         : std::nullopt,
       major_hypersparse_first,
       major_range_first,
@@ -652,6 +612,16 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
       major_value_range_start_offset);
   }
 
+  // FIXME: deprecated, replaced with copmute_number_of_edges (which works with or without edge
+  // masking)
+  edge_t number_of_edges() const
+  {
+    CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented.");
+    return this->number_of_edges_;
+  }
+
+  edge_t compute_number_of_edges(raft::handle_t const& handle) const;
+
   rmm::device_uvector<edge_t> compute_in_degrees(raft::handle_t const& handle) const;
   rmm::device_uvector<edge_t> compute_out_degrees(raft::handle_t const& handle) const;
 
@@ -751,15 +721,26 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
     return local_sorted_unique_edge_dst_vertex_partition_offsets_;
   }
 
+  void attach_edge_mask(edge_property_view_t<edge_t, uint32_t const*, bool> edge_mask_view)
+  {
+    edge_mask_view_ = edge_mask_view;
+  }
+
+  void clear_edge_mask() { edge_mask_view_ = std::nullopt; }
+
+  bool has_edge_mask() const { return edge_mask_view_.has_value(); }
+
+  std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view() const
+  {
+    return edge_mask_view_;
+  }
+
  private:
-  std::vector<edge_t const*> edge_partition_offsets_{};
-  std::vector<vertex_t const*> edge_partition_indices_{};
+  std::vector<raft::device_span<edge_t const>> edge_partition_offsets_{};
+  std::vector<raft::device_span<vertex_t const>> edge_partition_indices_{};
 
   // relevant only if we use the CSR + DCSR (or CSC + DCSC) hybrid format
-  std::optional<std::vector<vertex_t const*>> edge_partition_dcs_nzd_vertices_{};
-  std::optional<std::vector<vertex_t>> edge_partition_dcs_nzd_vertex_counts_{};
-
-  std::vector<edge_t> edge_partition_number_of_edges_{};
+  std::optional<std::vector<raft::device_span<vertex_t const>>> edge_partition_dcs_nzd_vertices_{};
 
   partition_t<vertex_t> partition_{};
 
@@ -796,6 +777,8 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
                      std::optional<raft::host_span<vertex_t const>>,
                      std::optional<std::byte> /* dummy */>
     local_sorted_unique_edge_dst_vertex_partition_offsets_{std::nullopt};
+
+  std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view_{std::nullopt};
 };
 
 // single-GPU version
@@ -808,9 +791,8 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
   static constexpr bool is_storage_transposed = store_transposed;
   static constexpr bool is_multi_gpu          = multi_gpu;
 
-  graph_view_t(raft::handle_t const& handle,
-               edge_t const* offsets,
-               vertex_t const* indices,
+  graph_view_t(raft::device_span<edge_t const> offsets,
+               raft::device_span<vertex_t const> indices,
                graph_view_meta_t<vertex_t, edge_t, store_transposed, multi_gpu> meta);
 
   std::vector<vertex_t> vertex_partition_range_offsets() const
@@ -924,11 +906,19 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
   {
     assert(partition_idx == 0);  // there is only one edge partition in single-GPU
     return edge_partition_view_t<vertex_t, edge_t, false>(
-      raft::device_span<edge_t const>(offsets_, offsets_ + (this->number_of_vertices() + 1)),
-      raft::device_span<vertex_t const>(indices_, indices_ + this->number_of_edges()),
-      this->number_of_vertices());
+      offsets_, indices_, this->number_of_vertices());
+  }
+
+  // FIXME: deprecated, replaced with copmute_number_of_edges (which works with or without edge
+  // masking)
+  edge_t number_of_edges() const
+  {
+    CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented.");
+    return this->number_of_edges_;
   }
 
+  edge_t compute_number_of_edges(raft::handle_t const& handle) const;
+
   rmm::device_uvector<edge_t> compute_in_degrees(raft::handle_t const& handle) const;
   rmm::device_uvector<edge_t> compute_out_degrees(raft::handle_t const& handle) const;
 
@@ -1016,12 +1006,28 @@ class graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu, std::enable_if
     return std::nullopt;
   }
 
+  void attach_edge_mask(edge_property_view_t<edge_t, uint32_t const*, bool> edge_mask_view)
+  {
+    edge_mask_view_ = edge_mask_view;
+  }
+
+  void clear_edge_mask() { edge_mask_view_ = std::nullopt; }
+
+  bool has_edge_mask() const { return edge_mask_view_.has_value(); }
+
+  std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view() const
+  {
+    return edge_mask_view_;
+  }
+
  private:
-  edge_t const* offsets_{nullptr};
-  vertex_t const* indices_{nullptr};
+  raft::device_span<edge_t const> offsets_{};
+  raft::device_span<vertex_t const> indices_{};
 
   // segment offsets based on vertex degree, relevant only if vertex IDs are renumbered
   std::optional<std::vector<vertex_t>> segment_offsets_{std::nullopt};
+
+  std::optional<edge_property_view_t<edge_t, uint32_t const*, bool>> edge_mask_view_{std::nullopt};
 };
 
 }  // namespace cugraph

cpp/include/cugraph/utilities/misc_utils.cuh

@@ -26,6 +26,8 @@
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/optional.h>
 
+#include <cuda/functional>
+
 #include <optional>
 #include <tuple>
 #include <vector>
@@ -44,7 +46,8 @@ std::tuple<std::vector<vertex_t>, std::vector<edge_t>> compute_offset_aligned_ed
 {
   auto search_offset_first = thrust::make_transform_iterator(
     thrust::make_counting_iterator(size_t{1}),
-    [approx_edge_chunk_size] __device__(auto i) { return i * approx_edge_chunk_size; });
+    cuda::proclaim_return_type<size_t>(
+      [approx_edge_chunk_size] __device__(auto i) { return i * approx_edge_chunk_size; }));
   auto num_chunks = (num_edges + approx_edge_chunk_size - 1) / approx_edge_chunk_size;
 
   if (num_chunks > 1) {