Merge branch 'branch-24.06' into pyg-dist

rapidsai · May 28, 2024 · b1c9b06 · b1c9b06
2 parents d2fff15 + 3156569
commit b1c9b06
Show file tree

Hide file tree

Showing 91 changed files with 5,136 additions and 1,898 deletions.
diff --git a/benchmarks/cugraph/standalone/bulk_sampling/cugraph_bulk_sampling.py b/benchmarks/cugraph/standalone/bulk_sampling/cugraph_bulk_sampling.py
@@ -344,7 +344,7 @@ def generate_rmat_dataset(
     del label_df
     gc.collect()
 
-    dask_label_df = dask_cudf.from_dask_dataframe(dask_label_df)
+    dask_label_df = dask_label_df.to_backend("cudf")
 
     node_offsets = {"paper": 0}
     edge_offsets = {("paper", "cites", "paper"): 0}

diff --git a/ci/test_python.sh b/ci/test_python.sh
@@ -3,10 +3,6 @@
 
 set -euo pipefail
 
-# TODO: Enable dask query planning (by default) once some bugs are fixed.
-# xref: https://github.com/rapidsai/cudf/issues/15027
-export DASK_DATAFRAME__QUERY_PLANNING=False
-
 # Support invoking test_python.sh outside the script directory
 cd "$(dirname "$(realpath "${BASH_SOURCE[0]}")")"/../
 

diff --git a/ci/test_wheel.sh b/ci/test_wheel.sh
@@ -3,10 +3,6 @@
 
 set -eoxu pipefail
 
-# TODO: Enable dask query planning (by default) once some bugs are fixed.
-# xref: https://github.com/rapidsai/cudf/issues/15027
-export DASK_DATAFRAME__QUERY_PLANNING=False
-
 package_name=$1
 package_dir=$2
 

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
@@ -180,6 +180,7 @@ set(CUGRAPH_SOURCES
     src/community/detail/refine_sg.cu
     src/community/detail/refine_mg.cu
     src/community/edge_triangle_count_sg.cu
+    src/community/edge_triangle_count_mg.cu
     src/community/detail/maximal_independent_moves_sg.cu
     src/community/detail/maximal_independent_moves_mg.cu
     src/detail/utility_wrappers.cu

diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp
@@ -2007,6 +2007,24 @@ void triangle_count(raft::handle_t const& handle,
                     raft::device_span<edge_t> counts,
                     bool do_expensive_check = false);
 
+/*
+ * @brief Compute edge triangle counts.
+ *
+ * Compute edge triangle counts for the entire set of edges.
+ *
+ * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type.
+ * @tparam edge_t Type of edge identifiers. Needs to be an integral type.
+ * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false)
+ * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+ * handles to various CUDA libraries) to run graph algorithms.
+ * @param graph_view Graph view object.
+ *
+ * @return edge_property_t containing the edge triangle count
+ */
+template <typename vertex_t, typename edge_t, bool multi_gpu>
+edge_property_t<graph_view_t<vertex_t, edge_t, false, multi_gpu>, edge_t> edge_triangle_count(
+  raft::handle_t const& handle, graph_view_t<vertex_t, edge_t, false, multi_gpu> const& graph_view);
+
 /*
  * @brief Compute K-Truss.
  *

diff --git a/cpp/include/cugraph/edge_partition_device_view.cuh b/cpp/include/cugraph/edge_partition_device_view.cuh
@@ -214,9 +214,9 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MajorIterator>
-  size_t compute_number_of_edges(MajorIterator major_first,
-                                 MajorIterator major_last,
-                                 rmm::cuda_stream_view stream) const
+  __host__ size_t compute_number_of_edges(MajorIterator major_first,
+                                          MajorIterator major_last,
+                                          rmm::cuda_stream_view stream) const
   {
     return dcs_nzd_vertices_ ? thrust::transform_reduce(
                                  rmm::exec_policy(stream),
@@ -250,7 +250,7 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
                                  thrust::plus<size_t>());
   }
 
-  rmm::device_uvector<edge_t> compute_local_degrees(rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees(rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(this->major_range_size(), stream);
     if (dcs_nzd_vertices_) {
@@ -277,9 +277,9 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MajorIterator>
-  rmm::device_uvector<edge_t> compute_local_degrees(MajorIterator major_first,
-                                                    MajorIterator major_last,
-                                                    rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees(MajorIterator major_first,
+                                                             MajorIterator major_last,
+                                                             rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(thrust::distance(major_first, major_last), stream);
     if (dcs_nzd_vertices_) {
@@ -306,10 +306,10 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MaskIterator, typename MajorIterator>
-  size_t compute_number_of_edges_with_mask(MaskIterator mask_first,
-                                           MajorIterator major_first,
-                                           MajorIterator major_last,
-                                           rmm::cuda_stream_view stream) const
+  __host__ size_t compute_number_of_edges_with_mask(MaskIterator mask_first,
+                                                    MajorIterator major_first,
+                                                    MajorIterator major_last,
+                                                    rmm::cuda_stream_view stream) const
   {
     return dcs_nzd_vertices_ ? thrust::transform_reduce(
                                  rmm::exec_policy(stream),
@@ -348,8 +348,8 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MaskIterator>
-  rmm::device_uvector<edge_t> compute_local_degrees_with_mask(MaskIterator mask_first,
-                                                              rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees_with_mask(
+    MaskIterator mask_first, rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(this->major_range_size(), stream);
     if (dcs_nzd_vertices_) {
@@ -384,10 +384,11 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MaskIterator, typename MajorIterator>
-  rmm::device_uvector<edge_t> compute_local_degrees_with_mask(MaskIterator mask_first,
-                                                              MajorIterator major_first,
-                                                              MajorIterator major_last,
-                                                              rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees_with_mask(
+    MaskIterator mask_first,
+    MajorIterator major_first,
+    MajorIterator major_last,
+    rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(thrust::distance(major_first, major_last), stream);
     if (dcs_nzd_vertices_) {
@@ -553,9 +554,9 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MajorIterator>
-  size_t compute_number_of_edges(MajorIterator major_first,
-                                 MajorIterator major_last,
-                                 rmm::cuda_stream_view stream) const
+  __host__ size_t compute_number_of_edges(MajorIterator major_first,
+                                          MajorIterator major_last,
+                                          rmm::cuda_stream_view stream) const
   {
     return thrust::transform_reduce(
       rmm::exec_policy(stream),
@@ -573,7 +574,7 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
       thrust::plus<size_t>());
   }
 
-  rmm::device_uvector<edge_t> compute_local_degrees(rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees(rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(this->major_range_size(), stream);
     thrust::transform(rmm::exec_policy(stream),
@@ -589,9 +590,9 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MajorIterator>
-  rmm::device_uvector<edge_t> compute_local_degrees(MajorIterator major_first,
-                                                    MajorIterator major_last,
-                                                    rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees(MajorIterator major_first,
+                                                             MajorIterator major_last,
+                                                             rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(thrust::distance(major_first, major_last), stream);
     thrust::transform(rmm::exec_policy(stream),
@@ -607,10 +608,10 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MaskIterator, typename MajorIterator>
-  size_t compute_number_of_edges_with_mask(MaskIterator mask_first,
-                                           MajorIterator major_first,
-                                           MajorIterator major_last,
-                                           rmm::cuda_stream_view stream) const
+  __host__ size_t compute_number_of_edges_with_mask(MaskIterator mask_first,
+                                                    MajorIterator major_first,
+                                                    MajorIterator major_last,
+                                                    rmm::cuda_stream_view stream) const
   {
     return thrust::transform_reduce(
       rmm::exec_policy(stream),
@@ -632,8 +633,8 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MaskIterator>
-  rmm::device_uvector<edge_t> compute_local_degrees_with_mask(MaskIterator mask_first,
-                                                              rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees_with_mask(
+    MaskIterator mask_first, rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(this->major_range_size(), stream);
     thrust::transform(
@@ -651,10 +652,11 @@ class edge_partition_device_view_t<vertex_t, edge_t, multi_gpu, std::enable_if_t
   }
 
   template <typename MaskIterator, typename MajorIterator>
-  rmm::device_uvector<edge_t> compute_local_degrees_with_mask(MaskIterator mask_first,
-                                                              MajorIterator major_first,
-                                                              MajorIterator major_last,
-                                                              rmm::cuda_stream_view stream) const
+  __host__ rmm::device_uvector<edge_t> compute_local_degrees_with_mask(
+    MaskIterator mask_first,
+    MajorIterator major_first,
+    MajorIterator major_last,
+    rmm::cuda_stream_view stream) const
   {
     rmm::device_uvector<edge_t> local_degrees(thrust::distance(major_first, major_last), stream);
     thrust::transform(