Skip to content

Commit

Permalink
renumbering implementation
Browse files Browse the repository at this point in the history
  • Loading branch information
ChuckHastings committed Jan 13, 2025
1 parent 03f328f commit 677942c
Show file tree
Hide file tree
Showing 5 changed files with 355 additions and 6 deletions.
1 change: 1 addition & 0 deletions cpp/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -541,6 +541,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)

Expand Down
12 changes: 6 additions & 6 deletions cpp/include/cugraph_c/graph_functions.h
Original file line number Diff line number Diff line change
Expand Up @@ -472,16 +472,16 @@ cugraph_error_code_t cugraph_decompress_to_edgelist(const cugraph_resource_handl
*
* @param [in] handle Handle for accessing resources
* @param [in] renumber_map Host array with the renumber map
* @param [in/out] src Device array of src vertices to renumber
* @param [in/out] dst 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
* @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* src,
cugraph_type_erased_device_array_view_t* dst,
cugraph_type_erased_device_array_view_t* srcs,
cugraph_type_erased_device_array_view_t* dsts,
cugraph_error_t** error);

#ifdef __cplusplus
Expand Down
215 changes: 215 additions & 0 deletions cpp/src/c_api/renumber_arbitrary_edgelist.cu
Original file line number Diff line number Diff line change
@@ -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;
}
2 changes: 2 additions & 0 deletions cpp/tests/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -885,6 +885,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. #
Expand Down
131 changes: 131 additions & 0 deletions cpp/tests/c_api/renumber_arbitrary_edgelist_test.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,131 @@
/*
* 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;
}

0 comments on commit 677942c

Please sign in to comment.