Merge branch 'investigate_leiden_inconsistencies' of github.com:chuck…

…hastings/cugraph into investigate_leiden_inconsistencies
rapidsai · Apr 18, 2024 · c877f0d · c877f0d
2 parents 2e060b1 + e5ad3ba
commit c877f0d
Show file tree

Hide file tree

Showing 13 changed files with 43 additions and 60 deletions.
diff --git a/ci/test_python.sh b/ci/test_python.sh
@@ -190,7 +190,7 @@ fi
 
 if [[ "${RAPIDS_CUDA_VERSION}" == "11.8.0" ]]; then
   if [[ "${RUNNER_ARCH}" != "ARM64" ]]; then
-    rapids-mamba-retry env create --force -f env.yaml -n test_cugraph_pyg
+    rapids-mamba-retry env create --yes -f env.yaml -n test_cugraph_pyg
 
     # Temporarily allow unbound variables for conda activation.
     set +u

diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml
@@ -25,9 +25,7 @@ dependencies:
 - doxygen
 - fsspec>=0.6.0
 - gcc_linux-64=11.*
-- gmock>=1.13.0
 - graphviz
-- gtest>=1.13.0
 - ipython
 - libcudf==24.6.*
 - libcugraphops==24.6.*

diff --git a/conda/environments/all_cuda-122_arch-x86_64.yaml b/conda/environments/all_cuda-122_arch-x86_64.yaml
@@ -27,9 +27,7 @@ dependencies:
 - doxygen
 - fsspec>=0.6.0
 - gcc_linux-64=11.*
-- gmock>=1.13.0
 - graphviz
-- gtest>=1.13.0
 - ipython
 - libcublas-dev
 - libcudf==24.6.*

diff --git a/conda/recipes/libcugraph/conda_build_config.yaml b/conda/recipes/libcugraph/conda_build_config.yaml
@@ -19,9 +19,6 @@ doxygen_version:
 nccl_version:
   - ">=2.9.9"
 
-gtest_version:
-  - ">=1.13.0"
-
 sysroot_version:
   - "2.17"
 

diff --git a/conda/recipes/libcugraph/meta.yaml b/conda/recipes/libcugraph/meta.yaml
@@ -68,8 +68,6 @@ requirements:
     {% endif %}
     - cuda-version ={{ cuda_version }}
     - doxygen {{ doxygen_version }}
-    - gmock {{ gtest_version }}
-    - gtest {{ gtest_version }}
     - libcudf ={{ minor_version }}
     - libcugraphops ={{ minor_version }}
     - libraft ={{ minor_version }}
@@ -190,8 +188,6 @@ outputs:
         {% else %}
         - cuda-cudart
         {% endif %}
-        - gmock {{ gtest_version }}
-        - gtest {{ gtest_version }}
     about:
       home: https://rapids.ai/
       dev_url: https://github.com/rapidsai/cugraph

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
@@ -151,7 +151,8 @@ if (BUILD_CUGRAPH_MTMG_TESTS)
 endif()
 
 if(BUILD_TESTS)
-  include(cmake/thirdparty/get_gtest.cmake)
+  include(${rapids-cmake-dir}/cpm/gtest.cmake)
+  rapids_cpm_gtest(BUILD_STATIC)
 endif()
 
 ################################################################################

diff --git a/cpp/cmake/thirdparty/get_gtest.cmake b/cpp/cmake/thirdparty/get_gtest.cmake
diff --git a/cpp/src/components/vertex_coloring_impl.cuh b/cpp/src/components/vertex_coloring_impl.cuh
@@ -89,21 +89,32 @@ rmm::device_uvector<vertex_t> vertex_coloring(
 
     if (current_graph_view.compute_number_of_edges(handle) == 0) { break; }
 
-    cugraph::edge_src_property_t<graph_view_t, flag_t> src_mis_flags(handle, current_graph_view);
-    cugraph::edge_dst_property_t<graph_view_t, flag_t> dst_mis_flags(handle, current_graph_view);
+    cugraph::edge_src_property_t<graph_view_t, flag_t> src_mis_flags(handle);
+    cugraph::edge_dst_property_t<graph_view_t, flag_t> dst_mis_flags(handle);
 
-    cugraph::update_edge_src_property(
-      handle, current_graph_view, is_vertex_in_mis.begin(), src_mis_flags);
+    if constexpr (graph_view_t::is_multi_gpu) {
+      src_mis_flags =
+        cugraph::edge_src_property_t<graph_view_t, flag_t>(handle, current_graph_view);
+      dst_mis_flags =
+        cugraph::edge_dst_property_t<graph_view_t, flag_t>(handle, current_graph_view);
 
-    cugraph::update_edge_dst_property(
-      handle, current_graph_view, is_vertex_in_mis.begin(), dst_mis_flags);
+      cugraph::update_edge_src_property(
+        handle, current_graph_view, is_vertex_in_mis.begin(), src_mis_flags);
+
+      cugraph::update_edge_dst_property(
+        handle, current_graph_view, is_vertex_in_mis.begin(), dst_mis_flags);
+    }
 
     if (color_id % 2 == 0) {
       cugraph::transform_e(
         handle,
         current_graph_view,
-        src_mis_flags.view(),
-        dst_mis_flags.view(),
+        graph_view_t::is_multi_gpu
+          ? src_mis_flags.view()
+          : detail::edge_major_property_view_t<vertex_t, flag_t const*>(is_vertex_in_mis.begin()),
+        graph_view_t::is_multi_gpu ? dst_mis_flags.view()
+                                   : detail::edge_minor_property_view_t<vertex_t, flag_t const*>(
+                                       is_vertex_in_mis.begin(), vertex_t{0}),
         cugraph::edge_dummy_property_t{}.view(),
         [color_id] __device__(
           auto src, auto dst, auto is_src_in_mis, auto is_dst_in_mis, thrust::nullopt_t) {
@@ -118,8 +129,12 @@ rmm::device_uvector<vertex_t> vertex_coloring(
       cugraph::transform_e(
         handle,
         current_graph_view,
-        src_mis_flags.view(),
-        dst_mis_flags.view(),
+        graph_view_t::is_multi_gpu
+          ? src_mis_flags.view()
+          : detail::edge_major_property_view_t<vertex_t, flag_t const*>(is_vertex_in_mis.begin()),
+        graph_view_t::is_multi_gpu ? dst_mis_flags.view()
+                                   : detail::edge_minor_property_view_t<vertex_t, flag_t const*>(
+                                       is_vertex_in_mis.begin(), vertex_t{0}),
         cugraph::edge_dummy_property_t{}.view(),
         [color_id] __device__(
           auto src, auto dst, auto is_src_in_mis, auto is_dst_in_mis, thrust::nullopt_t) {

diff --git a/cpp/src/prims/detail/extract_transform_v_frontier_e.cuh b/cpp/src/prims/detail/extract_transform_v_frontier_e.cuh
@@ -153,12 +153,15 @@ __global__ static void extract_transform_v_frontier_e_hypersparse_or_low_degree(
                                  typename EdgePartitionEdgeValueInputWrapper::value_type,
                                  EdgeOp>::type;
 
-  auto const tid          = threadIdx.x + blockIdx.x * blockDim.x;
-  auto const warp_id      = threadIdx.x / raft::warp_size();
-  auto const lane_id      = tid % raft::warp_size();
-  auto major_start_offset = static_cast<size_t>(*(edge_partition.major_hypersparse_first()) -
-                                                edge_partition.major_range_first());
-  auto idx                = static_cast<size_t>(tid);
+  auto const tid     = threadIdx.x + blockIdx.x * blockDim.x;
+  auto const warp_id = threadIdx.x / raft::warp_size();
+  auto const lane_id = tid % raft::warp_size();
+  [[maybe_unused]] vertex_t major_start_offset{};  // relevant only when hypersparse is true
+  if constexpr (hypersparse) {
+    major_start_offset = static_cast<size_t>(*(edge_partition.major_hypersparse_first()) -
+                                             edge_partition.major_range_first());
+  }
+  auto idx = static_cast<size_t>(tid);
 
   cuda::atomic_ref<size_t, cuda::thread_scope_device> buffer_idx(*buffer_idx_ptr);
 

diff --git a/cpp/src/prims/key_store.cuh b/cpp/src/prims/key_store.cuh
@@ -326,7 +326,7 @@ class key_cuco_store_t {
       rmm::mr::polymorphic_allocator<std::byte>(rmm::mr::get_current_device_resource()), stream);
     cuco_store_ =
       std::make_unique<cuco_set_type>(cuco_size,
-                                      cuco::sentinel::empty_key<key_t>{invalid_key},
+                                      cuco::empty_key<key_t>{invalid_key},
                                       thrust::equal_to<key_t>{},
                                       cuco::linear_probing<1,  // CG size
                                                            cuco::murmurhash3_32<key_t>>{},

diff --git a/cpp/src/prims/kv_store.cuh b/cpp/src/prims/kv_store.cuh
@@ -823,8 +823,8 @@ class kv_cuco_store_t {
     if constexpr (std::is_arithmetic_v<value_t>) {
       cuco_store_ =
         std::make_unique<cuco_map_type>(cuco_size,
-                                        cuco::sentinel::empty_key<key_t>{invalid_key},
-                                        cuco::sentinel::empty_value<value_t>{invalid_value},
+                                        cuco::empty_key<key_t>{invalid_key},
+                                        cuco::empty_value<value_t>{invalid_value},
                                         thrust::equal_to<key_t>{},
                                         cuco::linear_probing<1,  // CG size
                                                              cuco::murmurhash3_32<key_t>>{},
@@ -835,8 +835,8 @@ class kv_cuco_store_t {
     } else {
       cuco_store_ = std::make_unique<cuco_map_type>(
         cuco_size,
-        cuco::sentinel::empty_key<key_t>{invalid_key},
-        cuco::sentinel::empty_value<size_t>{std::numeric_limits<size_t>::max()},
+        cuco::empty_key<key_t>{invalid_key},
+        cuco::empty_value<size_t>{std::numeric_limits<size_t>::max()},
         thrust::equal_to<key_t>{},
         cuco::linear_probing<1,  // CG size
                              cuco::murmurhash3_32<key_t>>{},

diff --git a/cpp/tests/centrality/betweenness_centrality_reference.hpp b/cpp/tests/centrality/betweenness_centrality_reference.hpp
@@ -19,6 +19,7 @@
 #include <gtest/gtest.h>
 
 #include <limits>
+#include <optional>
 #include <queue>
 #include <stack>
 #include <vector>

diff --git a/dependencies.yaml b/dependencies.yaml
@@ -354,8 +354,6 @@ dependencies:
         packages:
           - c-compiler
           - cxx-compiler
-          - gmock>=1.13.0
-          - gtest>=1.13.0
           - libcudf==24.6.*
           - libcugraphops==24.6.*
           - libraft-headers==24.6.*