Merge branch 'branch-23.08' into sampling_updates_for_23.08

rapidsai · Jul 18, 2023 · ae962b1 · ae962b1
2 parents 7409c8e + 32e6e51
commit ae962b1
Show file tree

Hide file tree

Showing 13 changed files with 579 additions and 131 deletions.
diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml
@@ -30,6 +30,7 @@ jobs:
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
       date: ${{ inputs.date }}
@@ -39,6 +40,7 @@ jobs:
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: ${{ inputs.build_type || 'branch' }}
       branch: ${{ inputs.branch }}
       date: ${{ inputs.date }}
@@ -61,7 +63,7 @@ jobs:
       arch: "amd64"
       branch: ${{ inputs.branch }}
       build_type: ${{ inputs.build_type || 'branch' }}
-      container_image: "rapidsai/ci:latest"
+      container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10"
       date: ${{ inputs.date }}
       node_type: "gpu-v100-latest-1"
       run_script: "ci/build_docs.sh"

diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml
@@ -35,25 +35,29 @@ jobs:
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: pull-request
       node_type: cpu16
   conda-cpp-tests:
     needs: conda-cpp-build
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: pull-request
   conda-python-build:
     needs: conda-cpp-build
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: pull-request
   conda-python-tests:
     needs: conda-python-build
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: pull-request
   conda-notebook-tests:
     needs: conda-python-build
@@ -63,7 +67,7 @@ jobs:
       build_type: pull-request
       node_type: "gpu-v100-latest-1"
       arch: "amd64"
-      container_image: "rapidsai/ci:latest"
+      container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10"
       run_script: "ci/test_notebooks.sh"
   docs-build:
     needs: conda-python-build
@@ -73,7 +77,7 @@ jobs:
       build_type: pull-request
       node_type: "gpu-v100-latest-1"
       arch: "amd64"
-      container_image: "rapidsai/ci:latest"
+      container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10"
       run_script: "ci/build_docs.sh"
   wheel-build-pylibcugraph:
     needs: checks

diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml
@@ -18,6 +18,7 @@ jobs:
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: nightly
       branch: ${{ inputs.branch }}
       date: ${{ inputs.date }}
@@ -26,6 +27,7 @@ jobs:
     secrets: inherit
     uses: rapidsai/shared-action-workflows/.github/workflows/[email protected]
     with:
+      matrix_filter: map(select(.CUDA_VER | startswith("11")))
       build_type: nightly
       branch: ${{ inputs.branch }}
       date: ${{ inputs.date }}

diff --git a/cpp/include/cugraph/utilities/device_functors.cuh b/cpp/include/cugraph/utilities/device_functors.cuh
@@ -57,16 +57,28 @@ struct pack_bool_t {
   }
 };
 
-template <typename Iterator>
+template <typename index_t, typename Iterator>
 struct indirection_t {
   Iterator first{};
 
-  __device__ typename thrust::iterator_traits<Iterator>::value_type operator()(size_t i) const
+  __device__ typename thrust::iterator_traits<Iterator>::value_type operator()(index_t i) const
   {
     return *(first + i);
   }
 };
 
+template <typename index_t, typename Iterator>
+struct indirection_if_idx_valid_t {
+  Iterator first{};
+  index_t invalid_idx{};
+  typename thrust::iterator_traits<Iterator>::value_type invalid_value{};
+
+  __device__ typename thrust::iterator_traits<Iterator>::value_type operator()(index_t i) const
+  {
+    return (i != invalid_idx) ? *(first + i) : invalid_value;
+  }
+};
+
 template <typename T>
 struct not_equal_t {
   T compare{};

diff --git a/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh b/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh
@@ -27,7 +27,8 @@
 
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/device_atomics.cuh>
-#include <cudf/detail/utilities/hash_functions.cuh>
+#include <cudf/hashing/detail/default_hash.cuh>
+#include <cudf/hashing/detail/hash_functions.cuh>
 #include <cudf/utilities/error.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -118,7 +119,7 @@ union pair_packer<pair_type, std::enable_if_t<is_packable<pair_type>()>> {
  */
 template <typename Key,
           typename Element,
-          typename Hasher    = cudf::detail::default_hash<Key>,
+          typename Hasher    = cudf::hashing::detail::default_hash<Key>,
           typename Equality  = equal_to<Key>,
           typename Allocator = default_allocator<thrust::pair<Key, Element>>>
 class concurrent_unordered_map {

diff --git a/cpp/src/community/detail/common_methods.cuh b/cpp/src/community/detail/common_methods.cuh
@@ -18,6 +18,7 @@
 #include <community/detail/common_methods.hpp>
 
 #include <detail/graph_partition_utils.cuh>
+#include <prims/kv_store.cuh>
 #include <prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh>
 #include <prims/per_v_transform_reduce_incoming_outgoing_e.cuh>
 #include <prims/reduce_op.cuh>
@@ -42,6 +43,11 @@
 
 CUCO_DECLARE_BITWISE_COMPARABLE(float)
 CUCO_DECLARE_BITWISE_COMPARABLE(double)
+// FIXME: a temporary workaround for a compiler error, should be deleted once cuco gets patched.
+namespace cuco {
+template <>
+struct is_bitwise_comparable<cuco::pair<int32_t, float>> : std::true_type {};
+}  // namespace cuco
 
 namespace cugraph {
 namespace detail {

diff --git a/cpp/src/community/detail/refine_impl.cuh b/cpp/src/community/detail/refine_impl.cuh
@@ -48,6 +48,11 @@
 
 CUCO_DECLARE_BITWISE_COMPARABLE(float)
 CUCO_DECLARE_BITWISE_COMPARABLE(double)
+// FIXME: a temporary workaround for a compiler error, should be deleted once cuco gets patched.
+namespace cuco {
+template <>
+struct is_bitwise_comparable<cuco::pair<int32_t, float>> : std::true_type {};
+}  // namespace cuco
 
 namespace cugraph {
 namespace detail {

diff --git a/cpp/src/prims/detail/nbr_intersection.cuh b/cpp/src/prims/detail/nbr_intersection.cuh
@@ -974,7 +974,7 @@ nbr_intersection(raft::handle_t const& handle,
             .get_stream());  // initially store minimum degrees (upper bound for intersection sizes)
         if (intersect_minor_nbr[0] && intersect_minor_nbr[1]) {
           auto second_element_to_idx_map =
-            detail::kv_cuco_store_device_view_t((*major_to_idx_map_ptr)->view());
+            detail::kv_cuco_store_find_device_view_t((*major_to_idx_map_ptr)->view());
           thrust::transform(
             handle.get_thrust_policy(),
             get_dataframe_buffer_begin(vertex_pair_buffer),
@@ -1005,7 +1005,7 @@ nbr_intersection(raft::handle_t const& handle,
           handle.get_stream());
         if (intersect_minor_nbr[0] && intersect_minor_nbr[1]) {
           auto second_element_to_idx_map =
-            detail::kv_cuco_store_device_view_t((*major_to_idx_map_ptr)->view());
+            detail::kv_cuco_store_find_device_view_t((*major_to_idx_map_ptr)->view());
           thrust::tabulate(
             handle.get_thrust_policy(),
             rx_v_pair_nbr_intersection_sizes.begin(),