diff --git a/CHANGELOG.md b/CHANGELOG.md index 1a008279716..c96be8ca776 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,8 +1,8 @@ # cuGraph 0.8.0 (Date TBD) ## New Features -- PR #287 SNMG power iteration step1 -- PR #297 SNMG degree calculation +- PR #287 SNMG power iteration step1 +- PR #297 SNMG degree calculation - PR #300 Personalized Page Rank - PR #302 SNMG CSR Pagerank (cuda/C++) - PR #315 Weakly Connected Components adapted from cuML (cuda/C++) @@ -22,6 +22,7 @@ - PR #311 Fixed bug in SNMG degree causing failure for three gpus - PR #309 Update conda build recipes - PR #314 Added datasets to gitignore +- PR #322 Updates to accommodate new cudf include file locations # cuGraph 0.7.0 (10 May 2019) diff --git a/cpp/include/cugraph.h b/cpp/include/cugraph.h index d4531867bc8..12882fc9322 100644 --- a/cpp/include/cugraph.h +++ b/cpp/include/cugraph.h @@ -19,12 +19,10 @@ #include #include -#include +#include #include "types.h" #include "functions.h" #include "algorithms.h" - - diff --git a/cpp/src/converters/renumber.cuh b/cpp/src/converters/renumber.cuh index 5e2fa069267..6f7a3ccd251 100644 --- a/cpp/src/converters/renumber.cuh +++ b/cpp/src/converters/renumber.cuh @@ -30,7 +30,7 @@ #include #include -#include +#include #include #include "utilities/error_utils.h" @@ -44,7 +44,7 @@ namespace cugraph { typedef uint32_t hash_type; //typedef unsigned long long index_type; typedef uint32_t index_type; - + template class HashFunctionObject { public: @@ -236,13 +236,13 @@ namespace cugraph { [hash_bins_start, hash] __device__ (T_in vid) { atomicAdd(hash_bins_start + hash(vid), detail::index_type{1}); }); - + thrust::for_each(rmm::exec_policy(stream)->on(stream), dst, dst + size, [hash_bins_start, hash] __device__ (T_in vid) { atomicAdd(hash_bins_start + hash(vid), detail::index_type{1}); }); - + // // Compute exclusive sum and copy it into both hash_bins_start and @@ -261,7 +261,7 @@ namespace cugraph { detail::index_type hash_offset = atomicAdd(&hash_bins_end[hash_index], 1); hash_data[hash_offset] = vid; }); - + thrust::for_each(rmm::exec_policy(stream)->on(stream), dst, dst + size, [hash_bins_end, hash_data, hash] __device__ (T_in vid) { @@ -269,7 +269,7 @@ namespace cugraph { detail::index_type hash_offset = atomicAdd(&hash_bins_end[hash_index], 1); hash_data[hash_offset] = vid; }); - + // // Now we need to dedupe the hash bins // @@ -311,7 +311,7 @@ namespace cugraph { *new_size = temp; ALLOC_TRY(numbering_map, (*new_size) * sizeof(T_in), nullptr); - + T_in * local_numbering_map = *numbering_map; thrust::for_each(rmm::exec_policy(stream)->on(stream), @@ -332,7 +332,7 @@ namespace cugraph { return GDF_SUCCESS; } - + } #endif diff --git a/cpp/src/tests/components/con_comp_test.cu b/cpp/src/tests/components/con_comp_test.cu index 9cff418ac39..5fe823e4e8b 100755 --- a/cpp/src/tests/components/con_comp_test.cu +++ b/cpp/src/tests/components/con_comp_test.cu @@ -20,7 +20,7 @@ #include "gtest/gtest.h" #include "high_res_clock.h" -#include +#include #include "cuda_profiler_api.h" #include @@ -50,7 +50,7 @@ namespace{ //un-nammed { matrix_file = rhs.matrix_file; } - + Usecase& operator = (const Usecase& rhs) { matrix_file = rhs.matrix_file; @@ -64,19 +64,19 @@ namespace{ //un-nammed private: std::string matrix_file; }; - + }//end un-nammed namespace struct Tests_Weakly_CC : ::testing::TestWithParam { Tests_Weakly_CC() { } static void SetupTestCase() { } - static void TearDownTestCase() { + static void TearDownTestCase() { if (PERF) { for (unsigned int i = 0; i < weakly_cc_time.size(); ++i) { std::cout << weakly_cc_time[i] << std::endl; } - } + } } virtual void SetUp() { } virtual void TearDown() { } @@ -85,13 +85,13 @@ struct Tests_Weakly_CC : ::testing::TestWithParam void run_current_test(const Usecase& param) { const ::testing::TestInfo* const test_info =::testing::UnitTest::GetInstance()->current_test_info(); - std::stringstream ss; + std::stringstream ss; std::string test_id = std::string(test_info->test_case_name()) + std::string(".") + std::string(test_info->name()) + std::string("_") + getFileName(param.get_matrix_file())+ std::string("_") + ss.str().c_str(); cudaStream_t stream{nullptr}; int m, k, nnz; // MM_typecode mc; - + HighResClock hr_clock; double time_tmp; @@ -102,7 +102,7 @@ struct Tests_Weakly_CC : ::testing::TestWithParam ASSERT_TRUE(mm_is_matrix(mc)); ASSERT_TRUE(mm_is_coordinate(mc)); ASSERT_TRUE(mm_is_symmetric(mc));//weakly cc only works w/ undirected graphs, for now; - + // Allocate memory on host std::vector cooRowInd(nnz); std::vector cooColInd(nnz); @@ -144,10 +144,10 @@ struct Tests_Weakly_CC : ::testing::TestWithParam cudaDeviceSynchronize(); } EXPECT_EQ(status,GDF_SUCCESS); - + } }; - + std::vector Tests_Weakly_CC::weakly_cc_time; TEST_P(Tests_Weakly_CC, Weakly_CC) { @@ -155,7 +155,7 @@ TEST_P(Tests_Weakly_CC, Weakly_CC) { } // --gtest_filter=*simple_test* -INSTANTIATE_TEST_CASE_P(simple_test, Tests_Weakly_CC, +INSTANTIATE_TEST_CASE_P(simple_test, Tests_Weakly_CC, ::testing::Values( Usecase("networks/dolphins.mtx") //Usecase("networks/coPapersDBLP.mtx"), //Usecase("networks/coPapersCiteseer.mtx"), @@ -173,5 +173,3 @@ int main(int argc, char **argv) { return RUN_ALL_TESTS(); } - - diff --git a/cpp/src/utilities/grmat.cu b/cpp/src/utilities/grmat.cu index 8b5a50aacd7..4dd404e976e 100644 --- a/cpp/src/utilities/grmat.cu +++ b/cpp/src/utilities/grmat.cu @@ -32,7 +32,7 @@ #include -#include +#include #include #include "utilities/error_utils.h" #include "graph_utils.cuh" @@ -44,7 +44,7 @@ using namespace gunrock::util; using namespace gunrock::graphio; using namespace gunrock::graphio::grmat; -template +template __global__ void Remove_Self_Loops (VertexId* row, VertexId* col, Value* val, SizeT edges) { SizeT i = (SizeT)blockIdx.x * blockDim.x + threadIdx.x; @@ -54,7 +54,7 @@ __global__ void Remove_Self_Loops (VertexId* row, VertexId* col, Value* val, Siz if (row[i] == col[i]) { col[i] = 0; - } + } } } @@ -98,9 +98,9 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine bool quiet = false; typedef Coo_nv EdgeTupleType; - + cpu_timer.Start(); - + if (args->CheckCmdLineFlag ("rmat_scale") && args->CheckCmdLineFlag ("rmat_nodes")) { printf ("Please mention scale or nodes, not both \n"); @@ -111,8 +111,8 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine printf ("Please mention edgefactor or edge, not both \n"); return GDF_UNSUPPORTED_METHOD; } - - self_loops = args->CheckCmdLineFlag ("rmat_self_loops"); + + self_loops = args->CheckCmdLineFlag ("rmat_self_loops"); // graph construction or generation related parameters if (args -> CheckCmdLineFlag("normalized")) undirected = args -> CheckCmdLineFlag("rmat_undirected"); @@ -138,10 +138,10 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine rmat_seed = -1; } EdgeTupleType coo; - + if (undirected == true) { - rmat_all_edges = 2 * rmat_edges; + rmat_all_edges = 2 * rmat_edges; } else { @@ -169,8 +169,8 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine printf ("---------Graph properties-------\n" " Undirected : %s\n" " Nodes : %lld\n" - " Edges : %lld\n" - " a = %f, b = %f, c = %f, d = %f\n\n\n", ((undirected == true)? "True": "False"), (long long)rmat_nodes, + " Edges : %lld\n" + " a = %f, b = %f, c = %f, d = %f\n\n\n", ((undirected == true)? "True": "False"), (long long)rmat_nodes, (long long)(rmat_edges * ((undirected == true)? 2: 1)), rmat_a, rmat_b, rmat_c, rmat_d); } @@ -197,8 +197,8 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine cpu_timer2.Start(); cudaError_t status = cudaSuccess; if(val == nullptr) - status = BuildRmatGraph_coo_nv(rmat_nodes, rmat_edges, coo, undirected, - rmat_a, rmat_b, rmat_c, rmat_d, rmat_vmultipiler, rmat_vmin, rmat_seed, + status = BuildRmatGraph_coo_nv(rmat_nodes, rmat_edges, coo, undirected, + rmat_a, rmat_b, rmat_c, rmat_d, rmat_vmultipiler, rmat_vmin, rmat_seed, quiet, temp_devices.size(), gpu_idx); else status = BuildRmatGraph_coo_nv(rmat_nodes, rmat_edges, coo, undirected, @@ -222,15 +222,15 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine return GDF_CUDA_ERROR; } - + int block_size = (sizeof(VertexId) == 4) ? 1024 : 512; int grid_size = rmat_all_edges / block_size + 1; if (util::SetDevice(gpu_idx[0])) return GDF_CUDA_ERROR; if ((self_loops != false) && (val != nullptr)) - { - Remove_Self_Loops + { + Remove_Self_Loops <<>> (coo.row, coo.col, coo.val, rmat_all_edges); @@ -242,26 +242,26 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine VertexId nodes_row = 0; VertexId nodes_col = 0; - + cudaMemcpy((void*)&nodes_row, (void*)&(coo.row[rmat_all_edges-1]), sizeof(VertexId), cudaMemcpyDeviceToHost); - + tmp = thrust::max_element(rmm::exec_policy(stream)->on(stream), - thrust::device_pointer_cast((VertexId*)(coo.col)), + thrust::device_pointer_cast((VertexId*)(coo.col)), thrust::device_pointer_cast((VertexId*)(coo.col + rmat_all_edges))); nodes_col = tmp[0]; - + VertexId max_nodes = (nodes_row > nodes_col)? nodes_row: nodes_col; - + cpu_timer.Stop(); if ((src != nullptr) && (dest != nullptr)) - { - src->data = coo.row; + { + src->data = coo.row; src->size = rmat_all_edges; - src->valid = nullptr; - - dest->data = coo.col; - dest->size = rmat_all_edges; + src->valid = nullptr; + + dest->data = coo.col; + dest->size = rmat_all_edges; dest->valid = nullptr; } else @@ -274,20 +274,20 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine ALLOC_FREE_TRY(coo.val, stream); if (!quiet) printf ("Error : Pointers for gdf column are null, releasing allocated memory for graph\n"); - + return GDF_CUDA_ERROR; } - + if (val != nullptr) { - val->data = coo.val; - val->size = rmat_all_edges; + val->data = coo.val; + val->size = rmat_all_edges; val->valid = nullptr; } vertices = max_nodes+1; edges = rmat_all_edges; - + if (!quiet) printf ("Time to generate the graph %f ms\n" "Total time %f ms\n", cpu_timer2.ElapsedMillis(), cpu_timer.ElapsedMillis()); @@ -327,7 +327,7 @@ gdf_error gdf_grmat_gen (const char* argv, size_t& vertices, size_t& edges, gdf_ if (src == nullptr || dest == nullptr) { free_args(argc, arg); - return GDF_DATASET_EMPTY; + return GDF_DATASET_EMPTY; } GDF_REQUIRE ((src->dtype == dest->dtype), GDF_DTYPE_MISMATCH); diff --git a/python/cugraph/graph/c_graph.pxd b/python/cugraph/graph/c_graph.pxd index f164abc989e..5ce666e5d94 100755 --- a/python/cugraph/graph/c_graph.pxd +++ b/python/cugraph/graph/c_graph.pxd @@ -1,6 +1,6 @@ from libcpp cimport bool -cdef extern from "cudf.h": +cdef extern from "cudf/cudf.h": ctypedef size_t gdf_size_type @@ -16,7 +16,7 @@ cdef extern from "cudf.h": ctypedef struct gdf_dtype_extra_info: gdf_time_unit time_unit - ctypedef enum gdf_error: + ctypedef enum gdf_error: pass @@ -99,4 +99,4 @@ cdef extern from "cugraph.h": cdef gdf_error gdf_add_transposed_adj_list(gdf_graph *graph) cdef gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph) - cdef gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x) + cdef gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x)