From 71f143ee586f6d484ba5ffa30dcd14e6ad3bb7c4 Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Tue, 21 Jun 2022 20:36:53 -0400 Subject: [PATCH] Remove topology header (#2357) Remove the header file `topology.cuh` which is no longer used. Authors: - Chuck Hastings (https://github.com/ChuckHastings) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Rick Ratzel (https://github.com/rlratzel) URL: https://github.com/rapidsai/cugraph/pull/2357 --- cpp/src/components/connectivity.cu | 12 +- cpp/src/sampling/random_walks.cuh | 2 +- cpp/src/topology/topology.cuh | 397 ----------------------- cpp/tests/CMakeLists.txt | 4 - cpp/tests/components/scc_test.cu | 3 +- cpp/tests/sampling/rw_biased_seg_sort.cu | 141 -------- cpp/tests/sampling/rw_low_level_test.cu | 7 - 7 files changed, 8 insertions(+), 558 deletions(-) delete mode 100644 cpp/src/topology/topology.cuh delete mode 100644 cpp/tests/sampling/rw_biased_seg_sort.cu diff --git a/cpp/src/components/connectivity.cu b/cpp/src/components/connectivity.cu index 85134c1ad67..8668836dae5 100644 --- a/cpp/src/components/connectivity.cu +++ b/cpp/src/components/connectivity.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * Copyright (c) 2019-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -17,17 +17,17 @@ #include "scc_matrix.cuh" #include "weak_cc.cuh" -#include +#include -#include #include #include #include + +#include + +#include #include #include -#include - -#include "topology/topology.cuh" namespace cugraph { namespace detail { diff --git a/cpp/src/sampling/random_walks.cuh b/cpp/src/sampling/random_walks.cuh index dce5d1d4551..fe24d33f2b6 100644 --- a/cpp/src/sampling/random_walks.cuh +++ b/cpp/src/sampling/random_walks.cuh @@ -21,7 +21,6 @@ #include #include -#include #include #include @@ -29,6 +28,7 @@ #include +#include #include #include #include diff --git a/cpp/src/topology/topology.cuh b/cpp/src/topology/topology.cuh deleted file mode 100644 index fc86d219aba..00000000000 --- a/cpp/src/topology/topology.cuh +++ /dev/null @@ -1,397 +0,0 @@ -/* - * Copyright (c) 2019-2022, 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. - */ -#pragma once - -// Andrei Schaffer, 6/10/19, 7/23/21; -// -#include - -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include - -#include // for byte_t -#include -#include -#include -#include -#include - -namespace cugraph { -namespace topology { - -/** - * @brief Check symmetry of CSR adjacency matrix (raw pointers version); - * Algorithm outline: - * flag = true; - * for j in rows: - * for k in [row_offsets[j]..row_offsets[j+1]): - * col_indx = col_indices[k]; - * flag &= find(j, [col_indices[row_offsets[col_indx]]..col_indices[row_offsets[col_indx+1]]); - * return flag; - * - * @tparam IndexT type of indices for rows and columns - * @param handle raft handle - * @param nrows number of vertices - * @param ptr_r_o CSR row ofssets array - * @param nnz number of edges - * @param ptr_c_i CSR column indices array - */ -template -bool check_symmetry(raft::handle_t const& handle, - IndexT nrows, - const IndexT* ptr_r_o, - IndexT nnz, - const IndexT* ptr_c_i) -{ - using BoolT = bool; - rmm::device_uvector d_flags(nrows, handle.get_stream()); - thrust::fill_n(handle.get_thrust_policy(), d_flags.begin(), nrows, true); - - BoolT* start_flags = d_flags.data(); // d_flags.begin(); - BoolT* end_flags = start_flags + nrows; - BoolT init{1}; - return thrust::transform_reduce( - handle.get_thrust_policy(), - start_flags, - end_flags, - [ptr_r_o, ptr_c_i, start_flags, nnz] __device__(BoolT & crt_flag) { - IndexT row_indx = thrust::distance(start_flags, &crt_flag); - BoolT flag{1}; - for (auto k = ptr_r_o[row_indx]; k < ptr_r_o[row_indx + 1]; ++k) { - auto col_indx = ptr_c_i[k]; - auto begin = ptr_c_i + ptr_r_o[col_indx]; - auto end = - ptr_c_i + ptr_r_o[col_indx + 1]; // end is okay to point beyond last element of ptr_c_i - auto it = thrust::find(thrust::seq, begin, end, row_indx); - flag &= (it != end); - } - return crt_flag & flag; - }, - init, - thrust::logical_and()); -} - -#ifdef _DEBUG_ -/** - * @brief Sort weights of outgoing edges for each vertex; this requires a segmented (multi-partition - * parallel partial) sort, rather than complete sort. Required by Biased Random Walks. Caveat: this - * affects edge numbering. Naive implementation (slow) used only for debugging purposes, as it - * returns more information. - * - * Algorithm outline (version 1): - * input: num_vertices, num_edges, offsets[],indices[], weights[]; - * output: reordered indices[], weights[]; - * - * keys[] = sequence(num_edges); - * segment[] = vectorized-upper-bound(keys, offsets[]); // segment for each key, - * // to which key belongs - * sort-by_key(keys[], zip(indices[], weights[], [segment[], weights[]] (left_index, - * right_index){ - * // if left, right indices in the same segment then compare values: - * // - * if( segment[left_index] == segment[right_index] ) - * return weights[left_index] < weights[right_index]; - * else // "do nothing" - * return (left_index < right_index); // leave things unchanged - * }); - * - * The functor has on-demand instantiation semantics; meaning, the object is type agnostic - * (so that it can be instantiated by the caller without access to (ro, ci, vals)) - * while operator()(...) deduces its types from arguments and is meant to be called - * from inside `graph_t` memf's; - * operator()(...) returns std::tuple rather than thrust::tuple because the later has a bug with - * structured bindings; - */ -struct thrust_segment_sorter_by_weights_t { - thrust_segment_sorter_by_weights_t(raft::handle_t const& handle, size_t num_v, size_t num_e) - : handle_(handle), num_vertices_(num_v), num_edges_(num_e) - { - } - - template - std::tuple, rmm::device_uvector> operator()( - edge_t const* ptr_d_offsets_, vertex_t* ptr_d_indices_, weight_t* ptr_d_weights_) const - { - rmm::device_uvector d_keys(num_edges_, handle_.get_stream()); - rmm::device_uvector d_segs(num_edges_, handle_.get_stream()); - - // cannot use counting iterator, because d_keys gets passed to sort-by-key() - // - thrust::sequence(handle.get_thrust_policy(), d_keys.begin(), d_keys.end(), edge_t{0}); - - // d_segs = map each key(i.e., edge index), to corresponding - // segment (i.e., partition = out-going set) index - // - thrust::upper_bound(handle.get_thrust_policy(), - ptr_d_offsets_, - ptr_d_offsets_ + num_vertices_ + 1, - d_keys.begin(), - d_keys.end(), - d_segs.begin()); - - thrust::sort_by_key( - handle.get_thrust_policy(), - d_keys.begin(), - d_keys.end(), - thrust::make_zip_iterator(thrust::make_tuple(ptr_d_indices_, ptr_d_weights_)), - [ptr_segs = d_segs.data(), ptr_d_vals = ptr_d_weights_] __device__(auto left, auto right) { - // if both indices in same segment then compare the values: - // - if (ptr_segs[left] == ptr_segs[right]) { - return ptr_d_vals[left] < ptr_d_vals[right]; - } else { // don't compare them (leave things unchanged) - return (left < right); - } - }); - - return std::make_tuple(std::move(d_keys), std::move(d_segs)); - } - - private: - raft::handle_t const& handle_; - size_t num_vertices_; - size_t num_edges_; -}; -#endif - -/** - * @brief Sort weights of outgoing edges for each vertex; this requires a segmented (multi-partition - * parallel partial) sort, rather than complete sort. Required by Biased Random Walks. Caveat: this - * affects edge numbering. - * - * Algorithm outline (version 2): - * Uses cub::DeviceSegmentedRadixSort: - * https://nvlabs.github.io/cub/structcub_1_1_device_segmented_radix_sort.html - * - * The functor has on-demand instantiation semantics; meaning, the object is type agnostic - * (so that it can be instantiated by the caller without access to (ro, ci, vals)) - * while operator()(...) deduces its types from arguments and is meant to be called - * from inside `graph_t` memf's; - * It also must provide "in-place" semantics for modyfing the CSR array arguments; - */ -struct segment_sorter_by_weights_t { - segment_sorter_by_weights_t(raft::handle_t const& handle, size_t num_v, size_t num_e) - : handle_(handle), num_vertices_(num_v), num_edges_(num_e) - { - } - - template - std::tuple, rmm::device_uvector> operator()( - edge_t* ptr_d_offsets, vertex_t* ptr_d_indices, weight_t* ptr_d_weights) const - { - // keys are those on which sorting is done; hence, the weights: - // - rmm::device_uvector d_vals_out(num_edges_, handle_.get_stream()); - - // vals: are they just shuffled as a result of key sorting? - // no... it seems they participate in the sort... - // - rmm::device_uvector d_keys_out(num_edges_, handle_.get_stream()); - - // Note: In-place does not work; - - // Determine temporary device storage requirements: - // - void* ptr_d_temp_storage{nullptr}; - size_t temp_storage_bytes{0}; - cub::DeviceSegmentedRadixSort::SortPairs(ptr_d_temp_storage, - temp_storage_bytes, - ptr_d_weights, - d_vals_out.data(), // no in-place; - ptr_d_indices, - d_keys_out.data(), // no in-place; - num_edges_, - num_vertices_, - ptr_d_offsets, - ptr_d_offsets + 1, - 0, - (sizeof(weight_t) << 3), - handle_.get_stream()); - // Allocate temporary storage - // - rmm::device_uvector d_temp_storage(temp_storage_bytes, handle_.get_stream()); - ptr_d_temp_storage = d_temp_storage.data(); - - // Run sorting operation - // - cub::DeviceSegmentedRadixSort::SortPairs(ptr_d_temp_storage, - temp_storage_bytes, - ptr_d_weights, - d_vals_out.data(), // no in-place; - ptr_d_indices, - d_keys_out.data(), // no in-place; - num_edges_, - num_vertices_, - ptr_d_offsets, - ptr_d_offsets + 1, - 0, - (sizeof(weight_t) << 3), - handle_.get_stream()); - - // move data to deliver "in-place" semantics - // - return std::make_tuple(std::move(d_keys_out), std::move(d_vals_out)); - } - - template - void operator()(rmm::device_uvector& offsets, - rmm::device_uvector& indices, - std::optional>& weights) const - { - CUGRAPH_EXPECTS(weights.has_value(), "Cannot sort un-weighted graph by weights."); - - auto* ptr_d_offsets = offsets.data(); - auto* ptr_d_indices = indices.data(); - auto* ptr_d_weights = weights->data(); - - auto [d_keys_out, d_vals_out] = operator()(ptr_d_offsets, ptr_d_indices, ptr_d_weights); - - // move data to deliver "in-place" semantics - // - indices = std::move(d_keys_out); - *weights = std::move(d_vals_out); - } - - private: - raft::handle_t const& handle_; - size_t num_vertices_; - size_t num_edges_; -}; - -/** - * @brief Check if CSR's weights are segment-sorted, assuming a segment array is given; - * - * @tparam edge_t type of edge indices; - * @tparam edge_t type of weights; - * @param handle raft handle; - * @param ptr_d_segs array of segment index for each edge index; - * @param ptr_d_weights CSR weights array; - * @param num_edges number of edges; - */ -template -bool check_segmented_sort(raft::handle_t const& handle, - edge_t const* ptr_d_segs, - weight_t const* ptr_d_weights, - size_t num_edges) -{ - auto end = thrust::make_counting_iterator(num_edges - 1); - - // search for any adjacent elements in same segments - // that are _not_ ordered increasingly: - // - auto it = thrust::find_if( - handle.get_thrust_policy(), - thrust::make_counting_iterator(0), - end, - [ptr_d_segs, ptr_d_weights] __device__(auto indx) { - if (ptr_d_segs[indx] == ptr_d_segs[indx + 1]) { // consecutive indices in same segment - return (ptr_d_weights[indx] > ptr_d_weights[indx + 1]); - } else { // don't compare consecutive indices in different segments - return false; - } - }); - - return it == end; -} - -/** - * @brief Check if CSR's weights are segment-sorted, when no segment array is given; - * - * @tparam edge_t type of edge indices; - * @tparam edge_t type of weights; - * @param handle raft handle; - * @param ptr_d_offsets CSR offsets array; - * @param ptr_d_weights CSR weights array; - * @param num_vertices number of vertices; - * @param num_edges number of edges; - */ -template -bool check_segmented_sort(raft::handle_t const& handle, - edge_t const* ptr_d_offsets, - weight_t const* ptr_d_weights, - size_t num_vertices, - size_t num_edges) -{ - rmm::device_uvector d_keys(num_edges, handle.get_stream()); - rmm::device_uvector d_segs(num_edges, handle.get_stream()); - - // cannot use counting iterator, because d_keys gets passed to sort-by-key() - // - thrust::sequence(handle.get_thrust_policy(), d_keys.begin(), d_keys.end(), edge_t{0}); - - // d_segs = map each key(i.e., edge index), to corresponding - // segment (i.e., partition = out-going set) index - // - thrust::upper_bound(handle.get_thrust_policy(), - ptr_d_offsets, - ptr_d_offsets + num_vertices + 1, - d_keys.begin(), - d_keys.end(), - d_segs.begin()); - - return check_segmented_sort(handle, d_segs.data(), ptr_d_weights, num_edges); -} - -} // namespace topology -} // namespace cugraph - -namespace { // unnamed namespace for debugging tools: -template class Vector> -void print_v(const Vector& v, std::ostream& os) -{ - thrust::copy(v.begin(), v.end(), std::ostream_iterator(os, ",")); // okay - os << "\n"; -} - -template class Vector> -void print_v(const Vector& v, - typename Vector::const_iterator pos, - std::ostream& os) -{ - thrust::copy(v.begin(), pos, std::ostream_iterator(os, ",")); // okay - os << "\n"; -} - -template class Vector> -void print_v(const Vector& v, size_t n, std::ostream& os) -{ - thrust::copy_n(v.begin(), n, std::ostream_iterator(os, ",")); // okay - os << "\n"; -} - -template -void print_v(const T* p_v, size_t n, std::ostream& os) -{ - thrust::copy_n(p_v, n, std::ostream_iterator(os, ",")); // okay - os << "\n"; -} -} // namespace diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index fd370228a01..26fc60e2cd6 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -451,10 +451,6 @@ ConfigureTest(WEAKLY_CONNECTED_COMPONENTS_TEST components/weakly_connected_compo # - RANDOM_WALKS tests ---------------------------------------------------------------------------- ConfigureTest(RANDOM_WALKS_TEST sampling/random_walks_test.cu) -################################################################################################### -# - RANDOM_WALKS Segmented Sort tests ------------------------------------------------------------- -ConfigureTest(RANDOM_WALKS_SEG_SORT_TEST sampling/rw_biased_seg_sort.cu) - ################################################################################################### ConfigureTest(RANDOM_WALKS_LOW_LEVEL_TEST sampling/rw_low_level_test.cu) diff --git a/cpp/tests/components/scc_test.cu b/cpp/tests/components/scc_test.cu index beb78b21ba7..03dea413a08 100644 --- a/cpp/tests/components/scc_test.cu +++ b/cpp/tests/components/scc_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2019-2022, NVIDIA CORPORATION. All rights reserved. * * NVIDIA CORPORATION and its licensors retain all intellectual property * and proprietary rights in and to this software, related documentation @@ -23,7 +23,6 @@ #include #include #include -#include #include diff --git a/cpp/tests/sampling/rw_biased_seg_sort.cu b/cpp/tests/sampling/rw_biased_seg_sort.cu deleted file mode 100644 index 1f5a93211d1..00000000000 --- a/cpp/tests/sampling/rw_biased_seg_sort.cu +++ /dev/null @@ -1,141 +0,0 @@ -/* - * Copyright (c) 2021-2022, 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 "cuda_profiler_api.h" -#include "gtest/gtest.h" - -#include -#include - -#include - -#include - -#include - -#include - -#include -#include -#include -#include -#include -#include -#include - -namespace topo = cugraph::topology; - -struct RandomWalks_Usecase { - std::string graph_file_full_path{}; - bool test_weighted{false}; - - RandomWalks_Usecase(std::string const& graph_file_path, bool test_weighted) - : test_weighted(test_weighted) - { - if ((graph_file_path.length() > 0) && (graph_file_path[0] != '/')) { - graph_file_full_path = cugraph::test::get_rapids_dataset_root_dir() + "/" + graph_file_path; - } else { - graph_file_full_path = graph_file_path; - } - }; -}; - -class Tests_RWSegSort : public ::testing::TestWithParam { - public: - Tests_RWSegSort() {} - static void SetupTestCase() {} - static void TearDownTestCase() {} - - virtual void SetUp() {} - virtual void TearDown() {} - - template - void run_current_test(RandomWalks_Usecase const& target) - { - raft::handle_t handle{}; - - // debuf info: - // - // std::cout << "read graph file: " << configuration.graph_file_full_path << std::endl; - cugraph::graph_t graph(handle); - std::tie(graph, std::ignore) = - cugraph::test::read_graph_from_matrix_market_file( - handle, target.graph_file_full_path, target.test_weighted, false); - - size_t num_vertices = graph.number_of_vertices(); - size_t num_edges = graph.number_of_edges(); - - topo::segment_sorter_by_weights_t seg_sort(handle, num_vertices, num_edges); - - auto graph_view = graph.view(); - - // NOTE: barring a graph.sort() method, - // this const_cast<> is the only way to test - // segmented weight sort for a graph; - // - - edge_t* offsets = const_cast(graph_view.local_edge_partition_view().offsets()); - - vertex_t* indices = const_cast(graph_view.local_edge_partition_view().indices()); - weight_t* values = const_cast(*(graph_view.local_edge_partition_view().weights())); - - HighResTimer hr_timer; - std::string label{}; - - label = std::string("Biased RW: CUB Segmented Sort."); - hr_timer.start(label); - cudaProfilerStart(); - - auto [d_srt_indices, d_srt_weights] = seg_sort(offsets, indices, values); - - cudaProfilerStop(); - hr_timer.stop(); - - bool check_seg_sort = - topo::check_segmented_sort(handle, offsets, d_srt_weights.data(), num_vertices, num_edges); - ASSERT_TRUE(check_seg_sort); - - try { - auto runtime = hr_timer.get_average_runtime(label); - - std::cout << "Segmented Sort for Biased RW:\n"; - - } catch (std::exception const& ex) { - std::cerr << ex.what() << '\n'; - return; - - } catch (...) { - std::cerr << "ERROR: Unknown exception on timer label search." << '\n'; - return; - } - hr_timer.display(std::cout); - } -}; - -TEST_P(Tests_RWSegSort, Initialize_i32_i32_f) -{ - run_current_test(GetParam()); -} - -INSTANTIATE_TEST_SUITE_P( - simple_test, - Tests_RWSegSort, - ::testing::Values(RandomWalks_Usecase("test/datasets/karate.mtx", true), - RandomWalks_Usecase("test/datasets/web-Google.mtx", true), - RandomWalks_Usecase("test/datasets/ljournal-2008.mtx", true), - RandomWalks_Usecase("test/datasets/webbase-1M.mtx", true))); - -CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/sampling/rw_low_level_test.cu b/cpp/tests/sampling/rw_low_level_test.cu index 6aa72c32782..9c3e036dc38 100644 --- a/cpp/tests/sampling/rw_low_level_test.cu +++ b/cpp/tests/sampling/rw_low_level_test.cu @@ -17,7 +17,6 @@ #include "cuda_profiler_api.h" #include -#include #include #include @@ -1220,8 +1219,6 @@ TEST(RandomWalksUtility, PathsToCOO) TEST(BiasedRandomWalks, SelectorSmallGraph) { - namespace topo = cugraph::topology; - raft::handle_t handle{}; using vertex_t = int32_t; @@ -1336,8 +1333,6 @@ TEST(BiasedRandomWalks, SelectorSmallGraph) TEST(Node2VecRandomWalks, Node2VecSmallGraph) { - namespace topo = cugraph::topology; - raft::handle_t handle{}; using vertex_t = int32_t; @@ -1458,8 +1453,6 @@ TEST(Node2VecRandomWalks, Node2VecSmallGraph) TEST(Node2VecRandomWalks, CachedNode2VecSmallGraph) { - namespace topo = cugraph::topology; - raft::handle_t handle{}; using vertex_t = int32_t;