Merge pull request #4910 from rapidsai/branch-25.02

Forward-merge branch-25.02 into branch-25.04

cpp/CMakeLists.txt

@@ -548,6 +548,7 @@ add_library(cugraph_c
         src/c_api/allgather.cpp
         src/c_api/decompress_to_edgelist.cpp
         src/c_api/edgelist.cpp
+        src/c_api/renumber_arbitrary_edgelist.cu
         )
 add_library(cugraph::cugraph_c ALIAS cugraph_c)
 

cpp/include/cugraph_c/graph_functions.h

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2024, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2025, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -463,6 +463,27 @@ cugraph_error_code_t cugraph_decompress_to_edgelist(const cugraph_resource_handl
                                                     cugraph_edgelist_t** result,
                                                     cugraph_error_t** error);
 
+/**
+ * @brief     Renumber arbitrary edgelist
+ *
+ * This function is designed to assist renumbering graph vertices in the case where the
+ * the global vertex id list exceeds the GPU memory.  Renumbering is done in-place in the
+ * supplied @p src and @p dst parameters.
+ *
+ * @param [in]     handle         Handle for accessing resources
+ * @param [in]     renumber_map   Host array with the renumber map
+ * @param [in/out] srcs           Device array of src vertices to renumber
+ * @param [in/out] dsts           Device array of dst vertices to renumber
+ * @param [out] error             Pointer to an error object storing details of any error.  Will
+ *                                be populated if error code is not CUGRAPH_SUCCESS
+ */
+cugraph_error_code_t cugraph_renumber_arbitrary_edgelist(
+  const cugraph_resource_handle_t* handle,
+  const cugraph_type_erased_host_array_view_t* renumber_map,
+  cugraph_type_erased_device_array_view_t* srcs,
+  cugraph_type_erased_device_array_view_t* dsts,
+  cugraph_error_t** error);
+
 #ifdef __cplusplus
 }
 #endif

cpp/src/c_api/renumber_arbitrary_edgelist.cu

@@ -0,0 +1,190 @@
+/*
+ * Copyright (c) 2025, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "c_api/resource_handle.hpp"
+#include "c_api/utils.hpp"
+
+#include <cugraph_c/error.h>
+#include <cugraph_c/graph_functions.h>
+
+#include <cugraph/graph.hpp>
+#include <cugraph/utilities/error.hpp>
+
+#include <thrust/binary_search.h>
+#include <thrust/iterator/counting_iterator.h>
+
+namespace {
+
+template <typename vertex_t>
+cugraph_error_code_t renumber_arbitrary_edgelist(
+  raft::handle_t const& handle,
+  cugraph::c_api::cugraph_type_erased_host_array_view_t const* renumber_map,
+  cugraph::c_api::cugraph_type_erased_device_array_view_t* srcs,
+  cugraph::c_api::cugraph_type_erased_device_array_view_t* dsts)
+{
+  rmm::device_uvector<vertex_t> vertices(2 * srcs->size_, handle.get_stream());
+
+  thrust::copy_n(
+    handle.get_thrust_policy(), srcs->as_type<vertex_t>(), srcs->size_, vertices.data());
+  thrust::copy_n(handle.get_thrust_policy(),
+                 dsts->as_type<vertex_t>(),
+                 dsts->size_,
+                 vertices.data() + srcs->size_);
+
+  thrust::sort(handle.get_thrust_policy(), vertices.begin(), vertices.end());
+  vertices.resize(
+    thrust::distance(vertices.begin(),
+                     thrust::unique(handle.get_thrust_policy(), vertices.begin(), vertices.end())),
+    handle.get_stream());
+
+  vertices.shrink_to_fit(handle.get_stream());
+  rmm::device_uvector<vertex_t> ids(vertices.size(), handle.get_stream());
+  thrust::fill(handle.get_thrust_policy(),
+               ids.begin(),
+               ids.end(),
+               cugraph::invalid_vertex_id<vertex_t>::value);
+
+  raft::device_span<vertex_t const> vertices_span{vertices.data(), vertices.size()};
+  raft::device_span<vertex_t> ids_span{ids.data(), ids.size()};
+
+  // Read chunk of renumber_map in a loop, updating base offset to compute vertex id
+  // FIXME: Compute this as a function of free memory?  Or some value that keeps a
+  // particular GPU saturated?
+  size_t chunk_size = size_t{1} << 20;
+
+  rmm::device_uvector<vertex_t> renumber_chunk(chunk_size, handle.get_stream());
+
+  for (size_t chunk_base_offset = 0; chunk_base_offset < renumber_map->size_;
+       chunk_base_offset += chunk_size) {
+    size_t size = std::min(chunk_size, renumber_map->size_ - chunk_base_offset);
+    if (size < chunk_size) renumber_chunk.resize(size, handle.get_stream());
+
+    raft::update_device(renumber_chunk.data(),
+                        renumber_map->as_type<vertex_t>() + chunk_base_offset,
+                        size,
+                        handle.get_stream());
+
+    thrust::for_each(
+      handle.get_thrust_policy(),
+      thrust::make_counting_iterator<size_t>(0),
+      thrust::make_counting_iterator<size_t>(renumber_chunk.size()),
+      [chunk_base_offset,
+       renumber_chunk_span =
+         raft::device_span<vertex_t const>{renumber_chunk.data(), renumber_chunk.size()},
+       vertices_span,
+       ids_span] __device__(size_t idx) {
+        auto pos = thrust::lower_bound(
+          thrust::seq, vertices_span.begin(), vertices_span.end(), renumber_chunk_span[idx]);
+        if ((pos != vertices_span.end()) && (*pos == renumber_chunk_span[idx])) {
+          ids_span[thrust::distance(vertices_span.begin(), pos)] =
+            static_cast<vertex_t>(chunk_base_offset + idx);
+        }
+      });
+  }
+
+  CUGRAPH_EXPECTS(thrust::count(handle.get_thrust_policy(),
+                                ids.begin(),
+                                ids.end(),
+                                cugraph::invalid_vertex_id<vertex_t>::value) == 0,
+                  "some vertices were not renumbered");
+
+  thrust::transform(
+    handle.get_thrust_policy(),
+    srcs->as_type<vertex_t>(),
+    srcs->as_type<vertex_t>() + srcs->size_,
+    srcs->as_type<vertex_t>(),
+    [vertices_span, ids_span] __device__(vertex_t v) {
+      return ids_span[thrust::distance(
+        vertices_span.begin(),
+        thrust::lower_bound(thrust::seq, vertices_span.begin(), vertices_span.end(), v))];
+    });
+
+  thrust::transform(
+    handle.get_thrust_policy(),
+    dsts->as_type<vertex_t>(),
+    dsts->as_type<vertex_t>() + srcs->size_,
+    dsts->as_type<vertex_t>(),
+    [vertices_span, ids_span] __device__(vertex_t v) {
+      return ids_span[thrust::distance(
+        vertices_span.begin(),
+        thrust::lower_bound(thrust::seq, vertices_span.begin(), vertices_span.end(), v))];
+    });
+
+  return CUGRAPH_SUCCESS;
+}
+
+}  // namespace
+
+extern "C" cugraph_error_code_t cugraph_renumber_arbitrary_edgelist(
+  const cugraph_resource_handle_t* handle,
+  const cugraph_type_erased_host_array_view_t* renumber_map,
+  cugraph_type_erased_device_array_view_t* srcs,
+  cugraph_type_erased_device_array_view_t* dsts,
+  cugraph_error_t** error)
+{
+  cugraph::c_api::cugraph_type_erased_host_array_view_t const* h_renumber_map =
+    reinterpret_cast<cugraph::c_api::cugraph_type_erased_host_array_view_t const*>(renumber_map);
+  cugraph::c_api::cugraph_type_erased_device_array_view_t* d_srcs =
+    reinterpret_cast<cugraph::c_api::cugraph_type_erased_device_array_view_t*>(srcs);
+  cugraph::c_api::cugraph_type_erased_device_array_view_t* d_dsts =
+    reinterpret_cast<cugraph::c_api::cugraph_type_erased_device_array_view_t*>(dsts);
+
+  CAPI_EXPECTS(h_renumber_map->type_ == d_srcs->type_,
+               CUGRAPH_INVALID_INPUT,
+               "type of renumber map and src vertices must match",
+               *error);
+
+  CAPI_EXPECTS(h_renumber_map->type_ == d_dsts->type_,
+               CUGRAPH_INVALID_INPUT,
+               "type of renumber map and dst vertices must match",
+               *error);
+
+  CAPI_EXPECTS(
+    d_srcs->size_ == d_dsts->size_, CUGRAPH_INVALID_INPUT, "src and dst sizes must match", *error);
+
+  *error = nullptr;
+
+  try {
+    switch (h_renumber_map->type_) {
+      case cugraph_data_type_id_t::INT32: {
+        return renumber_arbitrary_edgelist<int32_t>(
+          *reinterpret_cast<cugraph::c_api::cugraph_resource_handle_t const*>(handle)->handle_,
+          h_renumber_map,
+          d_srcs,
+          d_dsts);
+      } break;
+      case cugraph_data_type_id_t::INT64: {
+        return renumber_arbitrary_edgelist<int64_t>(
+          *reinterpret_cast<cugraph::c_api::cugraph_resource_handle_t const*>(handle)->handle_,
+          h_renumber_map,
+          d_srcs,
+          d_dsts);
+      } break;
+      default: {
+        std::stringstream ss;
+        ss << "ERROR: Unsupported data type enum:" << static_cast<int>(h_renumber_map->type_);
+        *error =
+          reinterpret_cast<cugraph_error_t*>(new cugraph::c_api::cugraph_error_t{ss.str().c_str()});
+        return CUGRAPH_INVALID_INPUT;
+      }
+    }
+  } catch (std::exception const& ex) {
+    *error = reinterpret_cast<::cugraph_error_t*>(new cugraph::c_api::cugraph_error_t{ex.what()});
+    return CUGRAPH_UNKNOWN_ERROR;
+  }
+
+  return CUGRAPH_SUCCESS;
+}

cpp/tests/CMakeLists.txt

@@ -896,6 +896,8 @@ ConfigureCTest(CAPI_TRIANGLE_COUNT_TEST c_api/triangle_count_test.c)
 ConfigureCTest(CAPI_LOUVAIN_TEST c_api/louvain_test.c)
 ConfigureCTest(CAPI_LEIDEN_TEST c_api/leiden_test.c)
 ConfigureCTest(CAPI_ECG_TEST c_api/ecg_test.c)
+ConfigureCTest(CAPI_RENUMBER_ARBITRARY_EDGELIST_TEST c_api/renumber_arbitrary_edgelist_test.c)
+
 #############################################################################
 # Skipping due to CUDA 12.2 failure that traces back to RAFT                #
 # TODO: Uncomment this once the issue is fixed.                             #

cpp/tests/c_api/renumber_arbitrary_edgelist_test.c

@@ -0,0 +1,123 @@
+/*
+ * Copyright (c) 2025, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "c_test_utils.h" /* RUN_TEST */
+#include "cugraph_c/array.h"
+
+#include <cugraph_c/algorithms.h>
+#include <cugraph_c/graph.h>
+
+#include <math.h>
+
+typedef int32_t vertex_t;
+
+int generic_renumber_arbitrary_edgelist_test(vertex_t* h_src,
+                                             vertex_t* h_dst,
+                                             vertex_t* h_renumber_map,
+                                             size_t num_edges,
+                                             size_t renumber_map_size)
+{
+  int test_ret_value = 0;
+
+  cugraph_error_code_t ret_code = CUGRAPH_SUCCESS;
+  cugraph_error_t* ret_error;
+
+  cugraph_resource_handle_t* p_handle = NULL;
+
+  p_handle = cugraph_create_resource_handle(NULL);
+  TEST_ASSERT(test_ret_value, p_handle != NULL, "resource handle creation failed.");
+
+  cugraph_type_erased_device_array_t* srcs;
+  cugraph_type_erased_device_array_t* dsts;
+  cugraph_type_erased_device_array_view_t* srcs_view;
+  cugraph_type_erased_device_array_view_t* dsts_view;
+  cugraph_type_erased_host_array_view_t* renumber_map_view;
+
+  ret_code = cugraph_type_erased_device_array_create(p_handle, num_edges, INT32, &srcs, &ret_error);
+  TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "srcs create failed.");
+
+  ret_code = cugraph_type_erased_device_array_create(p_handle, num_edges, INT32, &dsts, &ret_error);
+  TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dsts create failed.");
+
+  srcs_view = cugraph_type_erased_device_array_view(srcs);
+  dsts_view = cugraph_type_erased_device_array_view(dsts);
+
+  ret_code = cugraph_type_erased_device_array_view_copy_from_host(
+    p_handle, srcs_view, (byte_t*)h_src, &ret_error);
+  TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed.");
+
+  ret_code = cugraph_type_erased_device_array_view_copy_from_host(
+    p_handle, dsts_view, (byte_t*)h_dst, &ret_error);
+  TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst copy_from_host failed.");
+
+  renumber_map_view =
+    cugraph_type_erased_host_array_view_create(h_renumber_map, renumber_map_size, INT32);
+
+  ret_code = cugraph_renumber_arbitrary_edgelist(
+    p_handle, renumber_map_view, srcs_view, dsts_view, &ret_error);
+
+  vertex_t h_renumbered_srcs[num_edges];
+  vertex_t h_renumbered_dsts[num_edges];
+
+  ret_code = cugraph_type_erased_device_array_view_copy_to_host(
+    p_handle, (byte_t*)h_renumbered_srcs, srcs_view, &ret_error);
+  TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed.");
+
+  ret_code = cugraph_type_erased_device_array_view_copy_to_host(
+    p_handle, (byte_t*)h_renumbered_dsts, dsts_view, &ret_error);
+  TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed.");
+
+  for (int i = 0; (i < num_edges) && (test_ret_value == 0); ++i) {
+    vertex_t renumbered_src = -1;
+    vertex_t renumbered_dst = -1;
+
+    for (size_t j = 0; (j < renumber_map_size) && ((renumbered_src < 0) || (renumbered_dst < 0));
+         ++j) {
+      if (h_src[i] == h_renumber_map[j]) renumbered_src = (vertex_t)j;
+      if (h_dst[i] == h_renumber_map[j]) renumbered_dst = (vertex_t)j;
+    }
+
+    TEST_ASSERT(test_ret_value, h_renumbered_srcs[i] == renumbered_src, "src results don't match");
+    TEST_ASSERT(test_ret_value, h_renumbered_dsts[i] == renumbered_dst, "dst results don't match");
+  }
+
+  cugraph_type_erased_device_array_free(dsts);
+  cugraph_type_erased_device_array_free(srcs);
+  cugraph_free_resource_handle(p_handle);
+  cugraph_error_free(ret_error);
+
+  return test_ret_value;
+}
+
+int test_renumbering()
+{
+  size_t num_edges         = 8;
+  size_t renumber_map_size = 6;
+
+  vertex_t h_src[]          = {0, 1, 1, 2, 2, 2, 3, 4};
+  vertex_t h_dst[]          = {1, 3, 4, 0, 1, 3, 5, 5};
+  vertex_t h_renumber_map[] = {5, 3, 1, 2, 4, 0};
+
+  return generic_renumber_arbitrary_edgelist_test(
+    h_src, h_dst, h_renumber_map, num_edges, renumber_map_size);
+}
+
+int main(int argc, char** argv)
+{
+  int result = 0;
+  result |= RUN_TEST(test_renumbering);
+  return result;
+}