diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index f1dda42ad9..0987757df2 100644 --- a/cpp/CMakeLists.txt +++ b/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) diff --git a/cpp/include/cugraph_c/graph_functions.h b/cpp/include/cugraph_c/graph_functions.h index 964b2f2c8d..4f9022888b 100644 --- a/cpp/include/cugraph_c/graph_functions.h +++ b/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 diff --git a/cpp/src/c_api/renumber_arbitrary_edgelist.cu b/cpp/src/c_api/renumber_arbitrary_edgelist.cu new file mode 100644 index 0000000000..644be0953e --- /dev/null +++ b/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 + +namespace { + +template +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 srcs_v(srcs->size_, handle.get_stream()); + rmm::device_uvector dsts_v(dsts->size_, handle.get_stream()); + rmm::device_uvector srcs_pos(srcs->size_, handle.get_stream()); + rmm::device_uvector dsts_pos(dsts->size_, handle.get_stream()); + + thrust::copy_n(handle.get_thrust_policy(), srcs->as_type(), srcs->size_, srcs_v.data()); + thrust::copy_n(handle.get_thrust_policy(), dsts->as_type(), 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 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() + chunk_base_offset, + size, + handle.get_stream()); + + rmm::device_uvector renumbered_values(srcs_v.size(), handle.get_stream()); + + thrust::fill(handle.get_thrust_policy(), + renumbered_values.begin(), + renumbered_values.end(), + cugraph::invalid_vertex_id::value); + + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(renumber_chunk.size()), + [chunk_base_offset, + renumber_chunk_span = + raft::device_span{renumber_chunk.data(), renumber_chunk.size()}, + srcs_span = raft::device_span{srcs_v.data(), srcs_v.size()}, + srcs_pos_span = raft::device_span{srcs_pos.data(), srcs_pos.size()}, + dsts_span = raft::device_span{dsts_v.data(), dsts_v.size()}, + dsts_pos_span = raft::device_span{dsts_pos.data(), dsts_pos.size()}, + output_srcs_span = raft::device_span{srcs->as_type(), srcs->size_}, + output_dsts_span = raft::device_span{dsts->as_type(), + dsts->size_}] __device__(size_t idx) { + vertex_t old_vertex_id = renumber_chunk_span[idx]; + vertex_t new_vertex_id = static_cast(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(); + ++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(); + ++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(); + })), + 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(); + })), + 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(renumber_map); + cugraph::c_api::cugraph_type_erased_device_array_view_t* d_srcs = + reinterpret_cast(srcs); + cugraph::c_api::cugraph_type_erased_device_array_view_t* d_dsts = + reinterpret_cast(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( + *reinterpret_cast(handle)->handle_, + h_renumber_map, + d_srcs, + d_dsts); + } break; + case cugraph_data_type_id_t::INT64: { + return renumber_arbitrary_edgelist( + *reinterpret_cast(handle)->handle_, + h_renumber_map, + d_srcs, + d_dsts); + } break; + default: { + std::stringstream ss; + ss << "ERROR: Unsupported data type enum:" << static_cast(h_renumber_map->type_); + *error = + reinterpret_cast(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; +} diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 6800b9c476..f819142076 100644 --- a/cpp/tests/CMakeLists.txt +++ b/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. # diff --git a/cpp/tests/c_api/renumber_arbitrary_edgelist_test.c b/cpp/tests/c_api/renumber_arbitrary_edgelist_test.c new file mode 100644 index 0000000000..6528ae5558 --- /dev/null +++ b/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 +#include + +#include + +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; +} diff --git a/python/pylibcugraph/pylibcugraph/CMakeLists.txt b/python/pylibcugraph/pylibcugraph/CMakeLists.txt index fe7c4b64aa..674f0769d6 100644 --- a/python/pylibcugraph/pylibcugraph/CMakeLists.txt +++ b/python/pylibcugraph/pylibcugraph/CMakeLists.txt @@ -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. You may obtain a copy of the License at @@ -71,6 +71,7 @@ set(cython_sources homogeneous_uniform_neighbor_sample.pyx edge_id_lookup_table.pyx decompress_to_edgelist.pyx + renumber_arbitrary_edgelist.pyx ) set(linked_libraries cugraph::cugraph;cugraph::cugraph_c) diff --git a/python/pylibcugraph/pylibcugraph/__init__.py b/python/pylibcugraph/pylibcugraph/__init__.py index f9561d961c..607b46bb6d 100644 --- a/python/pylibcugraph/pylibcugraph/__init__.py +++ b/python/pylibcugraph/pylibcugraph/__init__.py @@ -139,6 +139,8 @@ from pylibcugraph.decompress_to_edgelist import decompress_to_edgelist +from pylibcugraph.renumber_arbitrary_edgelist import renumber_arbitrary_edgelist + from pylibcugraph import exceptions diff --git a/python/pylibcugraph/pylibcugraph/_cugraph_c/graph_functions.pxd b/python/pylibcugraph/pylibcugraph/_cugraph_c/graph_functions.pxd index b27a7230a1..324ded80bb 100644 --- a/python/pylibcugraph/pylibcugraph/_cugraph_c/graph_functions.pxd +++ b/python/pylibcugraph/pylibcugraph/_cugraph_c/graph_functions.pxd @@ -1,4 +1,4 @@ -# 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. # You may obtain a copy of the License at @@ -34,7 +34,8 @@ from pylibcugraph._cugraph_c.similarity_algorithms cimport ( from pylibcugraph._cugraph_c.graph cimport cugraph_graph_t from pylibcugraph._cugraph_c.array cimport ( - cugraph_type_erased_device_array_view_t + cugraph_type_erased_device_array_view_t, + cugraph_type_erased_host_array_view_t, ) cdef extern from "cugraph_c/graph_functions.h": @@ -120,6 +121,14 @@ cdef extern from "cugraph_c/graph_functions.h": cugraph_error_t** error ) + cdef 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 + ) + ########################################################################### # induced_subgraph ctypedef struct cugraph_induced_subgraph_result_t: # Deprecated diff --git a/python/pylibcugraph/pylibcugraph/renumber_arbitrary_edgelist.pyx b/python/pylibcugraph/pylibcugraph/renumber_arbitrary_edgelist.pyx new file mode 100644 index 0000000000..067bc1cd3f --- /dev/null +++ b/python/pylibcugraph/pylibcugraph/renumber_arbitrary_edgelist.pyx @@ -0,0 +1,124 @@ +# 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. + +# Have cython use python 3 syntax +# cython: language_level = 3 + +from libc.stdint cimport uintptr_t + +from pylibcugraph.resource_handle cimport ResourceHandle + +from pylibcugraph._cugraph_c.resource_handle cimport ( + cugraph_resource_handle_t, +) +from pylibcugraph._cugraph_c.error cimport ( + cugraph_error_code_t, + cugraph_error_t, +) +from pylibcugraph._cugraph_c.array cimport ( + cugraph_type_erased_device_array_view_t, + cugraph_type_erased_host_array_view_t, + cugraph_type_erased_device_array_view_free, + cugraph_type_erased_device_array_view_create, + cugraph_type_erased_host_array_view_free, + cugraph_type_erased_host_array_view_create, +) + +from pylibcugraph._cugraph_c.graph_functions cimport ( + cugraph_renumber_arbitrary_edgelist, +) + +from pylibcugraph.utils cimport ( + assert_success, + assert_CAI_type, + assert_AI_type, + get_c_type_from_numpy_type, +) + +def renumber_arbitrary_edgelist( + ResourceHandle handle, + renumber_map, # host array + srcs, # device array + dsts, # device array +): + """ + Multi-GPU supporting function that accepts a local edgelist + and global renumber map and renumbers the edgelist in place. + + Parameters + ---------- + handle: ResourceHandle + Resource handle to use. + renumber_map: ndarray + Host array type containing the renumber map. + src: ndarray + Device array type containing the source vertices. + dst: ndarray + Device array type containing the destination vertices. + + Returns + ------- + Nothing. + """ + + assert_CAI_type(srcs, "srcs") + assert_CAI_type(dsts, "dsts") + + assert_AI_type(renumber_map, "renumber_map") + + cdef uintptr_t cai_renumber_map_ptr = \ + renumber_map.__array_interface__['data'][0] + cdef cugraph_type_erased_host_array_view_t* map_view = \ + cugraph_type_erased_host_array_view_create( + cai_renumber_map_ptr, + len(renumber_map), + get_c_type_from_numpy_type(renumber_map.dtype) + ) + + cdef uintptr_t cai_srcs_ptr = \ + srcs.__cuda_array_interface__['data'][0] + cdef cugraph_type_erased_device_array_view_t* srcs_view = \ + cugraph_type_erased_device_array_view_create( + cai_srcs_ptr, + len(srcs), + get_c_type_from_numpy_type(srcs.dtype) + ) + + cdef uintptr_t cai_dsts_ptr = \ + dsts.__cuda_array_interface__['data'][0] + cdef cugraph_type_erased_device_array_view_t* dsts_view = \ + cugraph_type_erased_device_array_view_create( + cai_dsts_ptr, + len(dsts), + get_c_type_from_numpy_type(dsts.dtype) + ) + + cdef cugraph_resource_handle_t* handle_cptr = handle.c_resource_handle_ptr + + cdef cugraph_error_t* err_cptr + + cdef cugraph_error_code_t err_code = cugraph_renumber_arbitrary_edgelist( + handle_cptr, + map_view, + srcs_view, + dsts_view, + &err_cptr, + ) + + # Verify that the C API call completed successfully and fail if it did not. + assert_success(err_code, err_cptr, "cugraph_renumber_arbitrary_edgelist") + + # Free the views + cugraph_type_erased_device_array_view_free(srcs_view) + cugraph_type_erased_device_array_view_free(dsts_view) + cugraph_type_erased_host_array_view_free(map_view) diff --git a/python/pylibcugraph/pylibcugraph/tests/test_renumber_arbitrary_edgelist.py b/python/pylibcugraph/pylibcugraph/tests/test_renumber_arbitrary_edgelist.py new file mode 100644 index 0000000000..1407dbdba7 --- /dev/null +++ b/python/pylibcugraph/pylibcugraph/tests/test_renumber_arbitrary_edgelist.py @@ -0,0 +1,34 @@ +# 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. + +import numpy as np +import cupy + +from pylibcugraph import renumber_arbitrary_edgelist +from pylibcugraph.resource_handle import ResourceHandle + + +def test_renumber_arbitrary_edgelist(): + renumber_map = np.array([5, 6, 1, 4, 0, 9]) + srcs = cupy.array([1, 1, 4, 4, 5, 5, 0, 9]) + dsts = cupy.array([6, 4, 5, 1, 4, 6, 1, 0]) + + renumber_arbitrary_edgelist( + ResourceHandle(), + renumber_map, + srcs, + dsts, + ) + + assert srcs.tolist() == [2, 2, 3, 3, 0, 0, 4, 5] + assert dsts.tolist() == [1, 3, 0, 2, 3, 1, 2, 4]