diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index b984c47c97b..febe4451882 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -168,10 +168,18 @@ ConfigureTest(PAGERANK_TEST "${PAGERANK_TEST_SRC}" "") # - SSSP tests ------------------------------------------------------------------------------------ set(SSSP_TEST_SRCS "${CMAKE_SOURCE_DIR}/../thirdparty/mmio/mmio.c" - "${CMAKE_CURRENT_SOURCE_DIR}/sssp/sssp_test.cu") + "${CMAKE_CURRENT_SOURCE_DIR}/traversal/sssp_test.cu") ConfigureTest(SSSP_TEST "${SSSP_TEST_SRCS}" "") +################################################################################################### +# - BFS tests ------------------------------------------------------------------------------------ +set(BFS_TEST_SRCS + "${CMAKE_SOURCE_DIR}/../thirdparty/mmio/mmio.c" + "${CMAKE_CURRENT_SOURCE_DIR}/traversal/bfs_test.cu") + +ConfigureTest(BFS_TEST "${BFS_TEST_SRCS}" "") + ################################################################################################### # - LOUVAIN tests --------------------------------------------------------------------------------- diff --git a/cpp/tests/centrality/betweenness_centrality_test.cu b/cpp/tests/centrality/betweenness_centrality_test.cu index b9870daf146..cdbb260c246 100644 --- a/cpp/tests/centrality/betweenness_centrality_test.cu +++ b/cpp/tests/centrality/betweenness_centrality_test.cu @@ -242,42 +242,6 @@ template void reference_betweenness_centrality( // ============================================================================= // Utility functions // ============================================================================= -// FIXME: This could be useful in other testsuite (SSSP, BFS, ...) -template -void generate_graph_csr(CSR_Result_Weighted &csr_result, - VT &m, - VT &nnz, - bool &is_directed, - std::string matrix_file) -{ - FILE *fpin = fopen(matrix_file.c_str(), "r"); - ASSERT_NE(fpin, nullptr) << "fopen (" << matrix_file << ") failure."; - - VT k; - MM_typecode mc; - ASSERT_EQ(mm_properties(fpin, 1, &mc, &m, &k, &nnz), 0) - << "could not read Matrix Market file properties" - << "\n"; - ASSERT_TRUE(mm_is_matrix(mc)); - ASSERT_TRUE(mm_is_coordinate(mc)); - ASSERT_FALSE(mm_is_complex(mc)); - ASSERT_FALSE(mm_is_skew(mc)); - is_directed = !mm_is_symmetric(mc); - - // Allocate memory on host - std::vector cooRowInd(nnz), cooColInd(nnz); - std::vector cooVal(nnz); - - // Read - ASSERT_EQ((mm_to_coo(fpin, 1, nnz, &cooRowInd[0], &cooColInd[0], &cooVal[0], NULL)), 0) - << "could not read matrix data" - << "\n"; - ASSERT_EQ(fclose(fpin), 0); - - ConvertCOOtoCSR_weighted(&cooRowInd[0], &cooColInd[0], &cooVal[0], nnz, csr_result); - CUDA_CHECK_LAST(); -} - // Compare while allowing relatie error of epsilon // zero_threshold indicates when we should drop comparison for small numbers template @@ -340,7 +304,8 @@ class Tests_BC : public ::testing::TestWithParam { ET nnz; CSR_Result_Weighted csr_result; bool is_directed = false; - generate_graph_csr(csr_result, m, nnz, is_directed, configuration.file_path_); + generate_graph_csr_from_mtx( + csr_result, m, nnz, is_directed, configuration.file_path_); cudaDeviceSynchronize(); cugraph::experimental::GraphCSRView G( csr_result.rowOffsets, csr_result.colIndices, csr_result.edgeWeights, m, nnz); @@ -446,7 +411,8 @@ class Tests_BFS : public ::testing::TestWithParam { ET nnz; CSR_Result_Weighted csr_result; bool is_directed = false; - generate_graph_csr(csr_result, m, nnz, is_directed, configuration.file_path_); + generate_graph_csr_from_mtx( + csr_result, m, nnz, is_directed, configuration.file_path_); cudaDeviceSynchronize(); cugraph::experimental::GraphCSRView G( csr_result.rowOffsets, csr_result.colIndices, csr_result.edgeWeights, m, nnz); @@ -517,8 +483,6 @@ class Tests_BFS : public ::testing::TestWithParam { //============================================================================== // Tests //============================================================================== -// BC -// ----------------------------------------------------------------------------- // Verifiy Un-Normalized results // Endpoint parameter is currently not usefull, is for later use TEST_P(Tests_BC, CheckFP32_NO_NORMALIZE_NO_ENDPOINTS) @@ -572,22 +536,6 @@ INSTANTIATE_TEST_CASE_P(simple_test, BC_Usecase("test/datasets/wiki2003.mtx", 4), BC_Usecase("test/datasets/wiki-Talk.mtx", 4))); -// BFS -// ----------------------------------------------------------------------------- -// FIXME: This could be reused by Issue #778 -TEST_P(Tests_BFS, CheckFP32) { run_current_test(GetParam()); } - -TEST_P(Tests_BFS, CheckFP64) { run_current_test(GetParam()); } - -INSTANTIATE_TEST_CASE_P(simple_test, - Tests_BFS, - ::testing::Values(BFS_Usecase("test/datasets/karate.mtx", 0), - BFS_Usecase("test/datasets/polbooks.mtx", 0), - BFS_Usecase("test/datasets/netscience.mtx", 0), - BFS_Usecase("test/datasets/netscience.mtx", 100), - BFS_Usecase("test/datasets/wiki2003.mtx", 1000), - BFS_Usecase("test/datasets/wiki-Talk.mtx", 1000))); - int main(int argc, char **argv) { rmmInitialize(nullptr); diff --git a/cpp/tests/test_utils.h b/cpp/tests/test_utils.h index 5794982ec0c..8da83880cc0 100644 --- a/cpp/tests/test_utils.h +++ b/cpp/tests/test_utils.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019, NVIDIA CORPORATION. + * Copyright (c) 2019-2020, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -47,9 +47,10 @@ extern "C" { #include #include "cugraph.h" - #include "utilities/error_utils.h" +#include "converters/COOtoCSR.cuh" + #ifndef CUDA_RT_CALL #define CUDA_RT_CALL(call) \ { \ @@ -143,7 +144,7 @@ void printv(size_t n, T* vec, int offset) dev_ptr + offset, dev_ptr + offset + n, std::ostream_iterator( - std::cout, " ")); // Assume no RMM dependency; TODO: check / test (potential BUG !!!!!) + std::cout, " ")); // Assume no RMM dependency; FIXME: check / test (potential BUG !!!!!) std::cout << std::endl; } @@ -841,8 +842,57 @@ bool gdf_csr_equal(gdf_column* a_off, gdf_column* a_ind, gdf_column* b_off, gdf_ return true; } +// ============================================================================ +// Generate data from an MTX, that can be used to Obtain a GraphCSRVIew +// ============================================================================ +// FIXME: A similar function could be useful for CSC format +// There are functions above that operate coo -> csr and coo->csc +/** + * @tparam + */ +template +void generate_graph_csr_from_mtx(CSR_Result_Weighted& csr_result, + VT& number_of_vertices, + VT& number_of_edges, + bool& directed, + std::string mtx_matrix_file) +{ + FILE* fpin = fopen(mtx_matrix_file.c_str(), "r"); + ASSERT_NE(fpin, nullptr) << "fopen (" << mtx_matrix_file << ") failure."; + + VT number_of_columns = 0; + MM_typecode mm_typecode{0}; + ASSERT_EQ(mm_properties( + fpin, 1, &mm_typecode, &number_of_vertices, &number_of_columns, &number_of_edges), + 0) + << "could not read Matrix Market file properties" + << "\n"; + ASSERT_TRUE(mm_is_matrix(mm_typecode)); + ASSERT_TRUE(mm_is_coordinate(mm_typecode)); + ASSERT_FALSE(mm_is_complex(mm_typecode)); + ASSERT_FALSE(mm_is_skew(mm_typecode)); + directed = !mm_is_symmetric(mm_typecode); + + // Allocate memory on host + std::vector coo_row_ind(number_of_edges); + std::vector coo_col_ind(number_of_edges); + std::vector coo_val(number_of_edges); + + // Read + ASSERT_EQ((mm_to_coo( + fpin, 1, number_of_edges, &coo_row_ind[0], &coo_col_ind[0], &coo_val[0], NULL)), + 0) + << "could not read matrix data" + << "\n"; + ASSERT_EQ(fclose(fpin), 0); + + ConvertCOOtoCSR_weighted( + &coo_row_ind[0], &coo_col_ind[0], &coo_val[0], number_of_edges, csr_result); + CUDA_CHECK_LAST(); +} + //////////////////////////////////////////////////////////////////////////////// -// TODO: move this code to rapids-core +// FIXME: move this code to rapids-core //////////////////////////////////////////////////////////////////////////////// // Define RAPIDS_DATASET_ROOT_DIR using a preprocessor variable to diff --git a/cpp/tests/traversal/bfs_test.cu b/cpp/tests/traversal/bfs_test.cu new file mode 100644 index 00000000000..da061f29529 --- /dev/null +++ b/cpp/tests/traversal/bfs_test.cu @@ -0,0 +1,281 @@ +/* + * Copyright (c) 2020, 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 governin_from_mtxg permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include + +#include "rmm_utils.h" + +#include "gtest/gtest.h" +#include "test_utils.h" + +// NOTE: This could be common to other files but we might not want the same precision +// depending on the algorithm +#ifndef TEST_EPSILON // It is currently use for relative error +#define TEST_EPSILON 0.0001 +#endif + +// NOTE: Defines under which values the difference should be discarded when +// considering values are close to zero +// i.e: Do we consider that the difference between 1.3e-9 and 8.e-12 is +// significant +#ifndef TEST_ZERO_THRESHOLD +#define TEST_ZERO_THRESHOLD 1e-10 +#endif +// ============================================================================ +// C++ Reference Implementation +// ============================================================================ +template +bool compare_close(const T &a, const T &b, const precision_t epsilon, precision_t zero_threshold) +{ + return ((zero_threshold > a && zero_threshold > b)) || + (a >= b * (1.0 - epsilon)) && (a <= b * (1.0 + epsilon)); +} +template +void populate_neighbors(VT *indices, ET *offsets, VT w, std::vector &neighbors) +{ + ET edge_start = offsets[w]; + ET edge_end = offsets[w + 1]; + ET edge_count = edge_end - edge_start; + + neighbors.clear(); // Reset neighbors vector's size + for (ET edge_idx = 0; edge_idx < edge_count; ++edge_idx) { + VT dst = indices[edge_start + edge_idx]; + neighbors.push_back(dst); + } +} + +// This implements the BFS from (Brandes, 2001) with shortest path counting +template +void ref_bfs(VT *indices, + ET *offsets, + VT const number_of_vertices, + std::queue &Q, + std::stack &S, + std::vector &dist, + std::vector> &pred, + std::vector &sigmas, + VT source) +{ + std::vector neighbors; + for (VT w = 0; w < number_of_vertices; ++w) { + pred[w].clear(); + dist[w] = std::numeric_limits::max(); + sigmas[w] = 0; + } + dist[source] = 0; + sigmas[source] = 1; + Q.push(source); + // b. Traversal + while (!Q.empty()) { + VT v = Q.front(); + Q.pop(); + S.push(v); + populate_neighbors(indices, offsets, v, neighbors); + for (VT w : neighbors) { + // Path Discovery: + // Found for the first time? + if (dist[w] == std::numeric_limits::max()) { + dist[w] = dist[v] + 1; + Q.push(w); + } + // Path counting + // Edge(v, w) on a shortest path? + if (dist[w] == dist[v] + 1) { + sigmas[w] += sigmas[v]; + pred[w].push_back(v); + } + } + } +} + +// ============================================================================ +// Test Suite +// ============================================================================ +typedef struct BFS_Usecase_t { + std::string config_; // Path to graph file + std::string file_path_; // Complete path to graph using dataset_root_dir + int source_; // Starting point from the traversal + BFS_Usecase_t(const std::string &config, int source) : config_(config), source_(source) + { + const std::string &rapidsDatasetRootDir = get_rapids_dataset_root_dir(); + if ((config_ != "") && (config_[0] != '/')) { + file_path_ = rapidsDatasetRootDir + "/" + config_; + } else { + file_path_ = config_; + } + }; +} BFS_Usecase; + +class Tests_BFS : public ::testing::TestWithParam { + public: + Tests_BFS() {} + static void SetupTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + template + void run_current_test(const BFS_Usecase &configuration) + { + // Step 1: Construction of the graph based on configuration + VT number_of_vertices; + ET number_of_edges; + CSR_Result_Weighted csr_result; + bool directed = false; + generate_graph_csr_from_mtx( + csr_result, number_of_vertices, number_of_edges, directed, configuration.file_path_); + CUDA_CHECK_LAST(); + cugraph::experimental::GraphCSRView G(csr_result.rowOffsets, + csr_result.colIndices, + csr_result.edgeWeights, + number_of_vertices, + number_of_edges); + CUDA_CHECK_LAST(); + G.prop.directed = directed; + + ASSERT_TRUE(configuration.source_ >= 0 && configuration.source_ <= G.number_of_vertices) + << "Starting sources should be >= 0 and" + << " less than the number of vertices in the graph"; + + VT source = configuration.source_; + + number_of_vertices = G.number_of_vertices; + number_of_edges = G.number_of_edges; + + std::vector indices(number_of_edges); + std::vector offsets(number_of_vertices + 1); + + CUDA_TRY( + cudaMemcpy(indices.data(), G.indices, sizeof(VT) * indices.size(), cudaMemcpyDeviceToHost)); + CUDA_TRY( + cudaMemcpy(offsets.data(), G.offsets, sizeof(ET) * offsets.size(), cudaMemcpyDeviceToHost)); + + std::queue Q; + std::stack S; + std::vector ref_bfs_dist(number_of_vertices); + std::vector> ref_bfs_pred(number_of_vertices); + std::vector ref_bfs_sigmas(number_of_vertices); + + ref_bfs(indices.data(), + offsets.data(), + number_of_vertices, + Q, + S, + ref_bfs_dist, + ref_bfs_pred, + ref_bfs_sigmas, + source); + + // Device data for cugraph_bfs + rmm::device_vector d_cugraph_dist(number_of_vertices); + rmm::device_vector d_cugraph_pred(number_of_vertices); + rmm::device_vector d_cugraph_sigmas(number_of_vertices); + + std::vector cugraph_dist(number_of_vertices); + std::vector cugraph_pred(number_of_vertices); + std::vector cugraph_sigmas(number_of_vertices); + + cugraph::bfs(G, + d_cugraph_dist.data().get(), + d_cugraph_pred.data().get(), + d_cugraph_sigmas.data().get(), + source, + G.prop.directed); + CUDA_TRY(cudaMemcpy(cugraph_dist.data(), + d_cugraph_dist.data().get(), + sizeof(VT) * d_cugraph_dist.size(), + cudaMemcpyDeviceToHost)); + CUDA_TRY(cudaMemcpy(cugraph_pred.data(), + d_cugraph_pred.data().get(), + sizeof(VT) * d_cugraph_pred.size(), + cudaMemcpyDeviceToHost)); + CUDA_TRY(cudaMemcpy(cugraph_sigmas.data(), + d_cugraph_sigmas.data().get(), + sizeof(double) * d_cugraph_sigmas.size(), + cudaMemcpyDeviceToHost)); + + for (VT i = 0; i < number_of_vertices; ++i) { + // Check distances: should be an exact match as we use signed int 32-bit + EXPECT_EQ(cugraph_dist[i], ref_bfs_dist[i]) + << "[MISMATCH] vaid = " << i << ", cugraph = " << cugraph_sigmas[i] + << " c++ ref = " << ref_bfs_sigmas[i]; + // Check predecessor: We do not enforce the predecessor, we simply verifiy + // that the predecessor obtained with the GPU implementation is one of the + // predecessors obtained during the C++ BFS traversal + VT pred = cugraph_pred[i]; // It could be equal to -1 if the node is never reached + if (pred == -1) { + EXPECT_TRUE(ref_bfs_pred[i].empty()) + << "[MISMATCH][PREDECESSOR] vaid = " << i << " cugraph had not predecessor," + << "while c++ ref found at least one."; + } else { + /* TODO(xcadet): Fix the following + std::vector::iterator it = + std::find(ref_bfs_pred[i].begin(), ref_bfs_pred[i].end(), pred); + EXPECT_TRUE(*it != ref_bfs_pred[i].end()) + << "[MISMATCH][PREDECESSOR] vaid = " << i << " cugraph = " << cugraph_sigmas[i] + << " , c++ ref did not consider it as a predecessor."; + */ + } + EXPECT_TRUE( + compare_close(cugraph_sigmas[i], ref_bfs_sigmas[i], TEST_EPSILON, TEST_ZERO_THRESHOLD)) + << "[MISMATCH] vaid = " << i << ", cugraph = " << cugraph_sigmas[i] + << " c++ ref = " << ref_bfs_sigmas[i]; + + if (return_sp_counter) { + EXPECT_TRUE( + compare_close(cugraph_sigmas[i], ref_bfs_sigmas[i], TEST_EPSILON, TEST_ZERO_THRESHOLD)) + << "[MISMATCH] vaid = " << i << ", cugraph = " << cugraph_sigmas[i] + << " c++ ref = " << ref_bfs_sigmas[i]; + } + } + } +}; + +//============================================================================== +// Tests +//============================================================================== +TEST_P(Tests_BFS, CheckFP32_NO_SP_COUNTER) { run_current_test(GetParam()); } + +TEST_P(Tests_BFS, CheckFP64_NO_SP_COUNTER) +{ + run_current_test(GetParam()); +} + +TEST_P(Tests_BFS, CheckFP32_SP_COUNTER) { run_current_test(GetParam()); } + +TEST_P(Tests_BFS, CheckFP64_SP_COUNTER) { run_current_test(GetParam()); } + +INSTANTIATE_TEST_CASE_P(simple_test, + Tests_BFS, + ::testing::Values(BFS_Usecase("test/datasets/karate.mtx", 0), + BFS_Usecase("test/datasets/polbooks.mtx", 0), + BFS_Usecase("test/datasets/netscience.mtx", 0), + BFS_Usecase("test/datasets/netscience.mtx", 100), + BFS_Usecase("test/datasets/wiki2003.mtx", 1000), + BFS_Usecase("test/datasets/wiki-Talk.mtx", 1000))); + +int main(int argc, char **argv) +{ + rmmInitialize(nullptr); + testing::InitGoogleTest(&argc, argv); + int rc = RUN_ALL_TESTS(); + rmmFinalize(); + return rc; +} diff --git a/cpp/tests/sssp/sssp_test.cu b/cpp/tests/traversal/sssp_test.cu similarity index 100% rename from cpp/tests/sssp/sssp_test.cu rename to cpp/tests/traversal/sssp_test.cu