Support C API Arbitrary Renumbering in PLC

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,215 @@
+/*
+ * 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/graph.hpp"
+#include "cugraph/utilities/error.hpp"
+#include "cugraph_c/error.h"
+#include "thrust/binary_search.h"
+#include "thrust/iterator/counting_iterator.h"
+#include "thrust/iterator/zip_iterator.h"
+
+#include <cugraph_c/graph_functions.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)
+{
+  // Create a sorted representation of each vertex id and where it exists in the input array
+  rmm::device_uvector<vertex_t> srcs_v(srcs->size_, handle.get_stream());
+  rmm::device_uvector<vertex_t> dsts_v(dsts->size_, handle.get_stream());
+  rmm::device_uvector<size_t> srcs_pos(srcs->size_, handle.get_stream());
+  rmm::device_uvector<size_t> dsts_pos(dsts->size_, handle.get_stream());
+
+  thrust::copy_n(handle.get_thrust_policy(), srcs->as_type<vertex_t>(), srcs->size_, srcs_v.data());
+  thrust::copy_n(handle.get_thrust_policy(), dsts->as_type<vertex_t>(), dsts->size_, dsts_v.data());
+  thrust::sequence(handle.get_thrust_policy(), srcs_pos.begin(), srcs_pos.end(), size_t{0});
+  thrust::sequence(handle.get_thrust_policy(), dsts_pos.begin(), dsts_pos.end(), size_t{0});
+
+  thrust::sort(handle.get_thrust_policy(),
+               thrust::make_zip_iterator(srcs_v.begin(), srcs_pos.begin()),
+               thrust::make_zip_iterator(srcs_v.end(), srcs_pos.end()));
+
+  thrust::sort(handle.get_thrust_policy(),
+               thrust::make_zip_iterator(dsts_v.begin(), dsts_pos.begin()),
+               thrust::make_zip_iterator(dsts_v.end(), dsts_pos.end()));
+
+  // 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());
+
+    rmm::device_uvector<vertex_t> renumbered_values(srcs_v.size(), handle.get_stream());
+
+    thrust::fill(handle.get_thrust_policy(),
+                 renumbered_values.begin(),
+                 renumbered_values.end(),
+                 cugraph::invalid_vertex_id<vertex_t>::value);
+
+    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()},
+       srcs_span        = raft::device_span<vertex_t>{srcs_v.data(), srcs_v.size()},
+       srcs_pos_span    = raft::device_span<size_t const>{srcs_pos.data(), srcs_pos.size()},
+       dsts_span        = raft::device_span<vertex_t>{dsts_v.data(), dsts_v.size()},
+       dsts_pos_span    = raft::device_span<size_t const>{dsts_pos.data(), dsts_pos.size()},
+       output_srcs_span = raft::device_span<vertex_t>{srcs->as_type<vertex_t>(), srcs->size_},
+       output_dsts_span = raft::device_span<vertex_t>{dsts->as_type<vertex_t>(),
+                                                      dsts->size_}] __device__(size_t idx) {
+        vertex_t old_vertex_id = renumber_chunk_span[idx];
+        vertex_t new_vertex_id = static_cast<vertex_t>(chunk_base_offset + idx);
+
+        auto begin_iter =
+          thrust::lower_bound(thrust::seq, srcs_span.begin(), srcs_span.end(), old_vertex_id);
+        if (begin_iter != srcs_span.end()) {
+          auto end_iter =
+            thrust::upper_bound(thrust::seq, srcs_span.begin(), srcs_span.end(), old_vertex_id);
+
+          while (begin_iter != end_iter) {
+            size_t offset = thrust::distance(srcs_span.begin(), begin_iter);
+            output_srcs_span[srcs_pos_span[offset]] = new_vertex_id;
+            srcs_span[offset]                       = cugraph::invalid_vertex_id<vertex_t>();
+            ++begin_iter;
+          }
+        }
+
+        begin_iter =
+          thrust::lower_bound(thrust::seq, dsts_span.begin(), dsts_span.end(), old_vertex_id);
+        if (begin_iter != dsts_span.end()) {
+          auto end_iter =
+            thrust::upper_bound(thrust::seq, dsts_span.begin(), dsts_span.end(), old_vertex_id);
+
+          while (begin_iter != end_iter) {
+            size_t offset = thrust::distance(dsts_span.begin(), begin_iter);
+            output_dsts_span[dsts_pos_span[offset]] = new_vertex_id;
+            dsts_span[offset]                       = cugraph::invalid_vertex_id<vertex_t>();
+            ++begin_iter;
+          }
+        }
+      });
+
+    srcs_v.resize(thrust::distance(
+                    thrust::make_zip_iterator(srcs_v.begin(), srcs_pos.begin()),
+                    thrust::remove_if(handle.get_thrust_policy(),
+                                      thrust::make_zip_iterator(srcs_v.begin(), srcs_pos.begin()),
+                                      thrust::make_zip_iterator(srcs_v.end(), srcs_pos.end()),
+                                      [] __device__(auto t) {
+                                        return thrust::get<0>(t) ==
+                                               cugraph::invalid_vertex_id<vertex_t>();
+                                      })),
+                  handle.get_stream());
+    srcs_pos.resize(srcs_v.size(), handle.get_stream());
+
+    dsts_v.resize(thrust::distance(
+                    thrust::make_zip_iterator(dsts_v.begin(), dsts_pos.begin()),
+                    thrust::remove_if(handle.get_thrust_policy(),
+                                      thrust::make_zip_iterator(dsts_v.begin(), dsts_pos.begin()),
+                                      thrust::make_zip_iterator(dsts_v.end(), dsts_pos.end()),
+                                      [] __device__(auto t) {
+                                        return thrust::get<0>(t) ==
+                                               cugraph::invalid_vertex_id<vertex_t>();
+                                      })),
+                  handle.get_stream());
+    dsts_pos.resize(dsts_v.size(), handle.get_stream());
+  }
+
+  CUGRAPH_EXPECTS(srcs_v.size() == 0, "some src vertices were not renumbered");
+  CUGRAPH_EXPECTS(dsts_v.size() == 0, "some dst vertices were not renumbered");
+
+  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);
+
+  *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

@@ -893,6 +893,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;
+}