From 0e2dff6dd65edbc6de346ff278400e2df8aba94c Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Tue, 7 May 2019 09:25:33 -0700 Subject: [PATCH 01/16] Fix file permissions --- cpp/src/bfs.cuh | 0 1 file changed, 0 insertions(+), 0 deletions(-) mode change 100755 => 100644 cpp/src/bfs.cuh diff --git a/cpp/src/bfs.cuh b/cpp/src/bfs.cuh old mode 100755 new mode 100644 From 4213a93bd6db43c35040839b5a5ea43ff17e1e60 Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Wed, 8 May 2019 10:38:59 -0700 Subject: [PATCH 02/16] Moved structure functions --- cpp/CMakeLists.txt | 1 + cpp/src/cugraph.cu | 273 ++-------------------------------- cpp/src/structure/cugraph.cu | 276 +++++++++++++++++++++++++++++++++++ 3 files changed, 287 insertions(+), 263 deletions(-) create mode 100644 cpp/src/structure/cugraph.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 464fed08342..2de2f5d753a 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -237,6 +237,7 @@ add_library(cugraph SHARED src/overlap.cu src/nvgraph_gdf.cu src/two_hop_neighbors.cu + src/structure/cugraph.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/test_utils.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/error_utils.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/misc_utils.cu diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu index 80975930de7..b277720c96a 100644 --- a/cpp/src/cugraph.cu +++ b/cpp/src/cugraph.cu @@ -64,60 +64,6 @@ void gdf_col_release(gdf_column* col) { delete col; } -void cpy_column_view(const gdf_column *in, gdf_column *out) { - if (in != nullptr && out !=nullptr) { - gdf_column_view(out, in->data, in->valid, in->size, in->dtype); - } -} - -gdf_error gdf_adj_list_view(gdf_graph *graph, const gdf_column *offsets, - const gdf_column *indices, const gdf_column *edge_data) { - //This function returns an error if this graph object has at least one graph - //representation to prevent a single object storing two different graphs. - GDF_REQUIRE( ((graph->edgeList == nullptr) && (graph->adjList == nullptr) && - (graph->transposedAdjList == nullptr)), GDF_INVALID_API_CALL); - GDF_REQUIRE( offsets->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( (offsets->dtype == indices->dtype), GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( ((offsets->dtype == GDF_INT32) || (offsets->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( (offsets->size > 0), GDF_DATASET_EMPTY ); - - graph->adjList = new gdf_adj_list; - graph->adjList->offsets = new gdf_column; - graph->adjList->indices = new gdf_column; - graph->adjList->ownership = 0; - - cpy_column_view(offsets, graph->adjList->offsets); - cpy_column_view(indices, graph->adjList->indices); - if (edge_data) { - GDF_REQUIRE( indices->size == edge_data->size, GDF_COLUMN_SIZE_MISMATCH ); - graph->adjList->edge_data = new gdf_column; - cpy_column_view(edge_data, graph->adjList->edge_data); - } - else { - graph->adjList->edge_data = nullptr; - } - return GDF_SUCCESS; -} - -gdf_error gdf_adj_list::get_vertex_identifiers(gdf_column *identifiers) { - GDF_REQUIRE( offsets != nullptr , GDF_INVALID_API_CALL); - GDF_REQUIRE( offsets->data != nullptr , GDF_INVALID_API_CALL); - cugraph::sequence((int)offsets->size-1, (int*)identifiers->data); - return GDF_SUCCESS; -} - -gdf_error gdf_adj_list::get_source_indices (gdf_column *src_indices) { - GDF_REQUIRE( offsets != nullptr , GDF_INVALID_API_CALL); - GDF_REQUIRE( offsets->data != nullptr , GDF_INVALID_API_CALL); - GDF_REQUIRE( src_indices->size == indices->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( src_indices->dtype == indices->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( src_indices->size > 0, GDF_DATASET_EMPTY ); - cugraph::offsets_to_indices((int*)offsets->data, offsets->size-1, (int*)src_indices->data); - - return GDF_SUCCESS; -} - gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, gdf_column *src_renumbered, gdf_column *dst_renumbered, gdf_column *numbering_map) { @@ -125,7 +71,7 @@ gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, GDF_REQUIRE( src->size == dst->size, GDF_COLUMN_SIZE_MISMATCH ); GDF_REQUIRE( src->dtype == dst->dtype, GDF_UNSUPPORTED_DTYPE ); GDF_REQUIRE( ((src->dtype == GDF_INT32) || (src->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( src->size > 0, GDF_DATASET_EMPTY ); + GDF_REQUIRE( src->size > 0, GDF_DATASET_EMPTY ); // // TODO: we're currently renumbering without using valid. We need to @@ -153,7 +99,7 @@ gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); gdf_column_view(src_renumbered, tmp, src->valid, src->size, src->dtype); - + ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, dst->dtype); @@ -182,7 +128,7 @@ gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, int64_t *tmp; ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); gdf_column_view(src_renumbered, tmp, src->valid, src->size, GDF_INT32); - + ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, GDF_INT32); @@ -213,143 +159,6 @@ gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, return GDF_SUCCESS; } -gdf_error gdf_edge_list_view(gdf_graph *graph, const gdf_column *src_indices, - const gdf_column *dest_indices, const gdf_column *edge_data) { - //This function returns an error if this graph object has at least one graph - //representation to prevent a single object storing two different graphs. - GDF_REQUIRE( ((graph->edgeList == nullptr) && (graph->adjList == nullptr) && - (graph->transposedAdjList == nullptr)), GDF_INVALID_API_CALL); - GDF_REQUIRE( src_indices->size == dest_indices->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( src_indices->dtype == dest_indices->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( ((src_indices->dtype == GDF_INT32) || (src_indices->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( src_indices->size > 0, GDF_DATASET_EMPTY ); - GDF_REQUIRE( src_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( dest_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - - graph->edgeList = new gdf_edge_list; - graph->edgeList->src_indices = new gdf_column; - graph->edgeList->dest_indices = new gdf_column; - graph->edgeList->ownership = 0; - - cpy_column_view(src_indices, graph->edgeList->src_indices); - cpy_column_view(dest_indices, graph->edgeList->dest_indices); - if (edge_data) { - GDF_REQUIRE( src_indices->size == edge_data->size, GDF_COLUMN_SIZE_MISMATCH ); - graph->edgeList->edge_data = new gdf_column; - cpy_column_view(edge_data, graph->edgeList->edge_data); - } - else { - graph->edgeList->edge_data = nullptr; - } - - return GDF_SUCCESS; -} - -template -gdf_error gdf_add_adj_list_impl (gdf_graph *graph) { - if (graph->adjList == nullptr) { - GDF_REQUIRE( graph->edgeList != nullptr , GDF_INVALID_API_CALL); - int nnz = graph->edgeList->src_indices->size, status = 0; - graph->adjList = new gdf_adj_list; - graph->adjList->offsets = new gdf_column; - graph->adjList->indices = new gdf_column; - graph->adjList->ownership = 1; - - if (graph->edgeList->edge_data!= nullptr) { - graph->adjList->edge_data = new gdf_column; - - CSR_Result_Weighted adj_list; - status = ConvertCOOtoCSR_weighted((int*)graph->edgeList->src_indices->data, (int*)graph->edgeList->dest_indices->data, (WT*)graph->edgeList->edge_data->data, nnz, adj_list); - - gdf_column_view(graph->adjList->offsets, adj_list.rowOffsets, - nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); - gdf_column_view(graph->adjList->indices, adj_list.colIndices, - nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); - gdf_column_view(graph->adjList->edge_data, adj_list.edgeWeights, - nullptr, adj_list.nnz, graph->edgeList->edge_data->dtype); - } - else { - CSR_Result adj_list; - status = ConvertCOOtoCSR((int*)graph->edgeList->src_indices->data,(int*)graph->edgeList->dest_indices->data, nnz, adj_list); - gdf_column_view(graph->adjList->offsets, adj_list.rowOffsets, - nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); - gdf_column_view(graph->adjList->indices, adj_list.colIndices, - nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); - } - if (status !=0) { - std::cerr << "Could not generate the adj_list" << std::endl; - return GDF_CUDA_ERROR; - } - } - return GDF_SUCCESS; -} - -gdf_error gdf_add_edge_list (gdf_graph *graph) { - if (graph->edgeList == nullptr) { - GDF_REQUIRE( graph->adjList != nullptr , GDF_INVALID_API_CALL); - int *d_src; - graph->edgeList = new gdf_edge_list; - graph->edgeList->src_indices = new gdf_column; - graph->edgeList->dest_indices = new gdf_column; - graph->edgeList->ownership = 2; - - CUDA_TRY(cudaMallocManaged ((void**)&d_src, sizeof(int) * graph->adjList->indices->size)); - - cugraph::offsets_to_indices((int*)graph->adjList->offsets->data, - graph->adjList->offsets->size-1, - (int*)d_src); - - gdf_column_view(graph->edgeList->src_indices, d_src, - nullptr, graph->adjList->indices->size, graph->adjList->indices->dtype); - cpy_column_view(graph->adjList->indices, graph->edgeList->dest_indices); - - if (graph->adjList->edge_data != nullptr) { - graph->edgeList->edge_data = new gdf_column; - cpy_column_view(graph->adjList->edge_data, graph->edgeList->edge_data); - } - } - return GDF_SUCCESS; -} - - -template -gdf_error gdf_add_transposed_adj_list_impl (gdf_graph *graph) { - if (graph->transposedAdjList == nullptr ) { - GDF_REQUIRE( graph->edgeList != nullptr , GDF_INVALID_API_CALL); - int nnz = graph->edgeList->src_indices->size, status = 0; - graph->transposedAdjList = new gdf_adj_list; - graph->transposedAdjList->offsets = new gdf_column; - graph->transposedAdjList->indices = new gdf_column; - graph->transposedAdjList->ownership = 1; - - if (graph->edgeList->edge_data) { - graph->transposedAdjList->edge_data = new gdf_column; - CSR_Result_Weighted adj_list; - status = ConvertCOOtoCSR_weighted( (int*)graph->edgeList->dest_indices->data, (int*)graph->edgeList->src_indices->data, (WT*)graph->edgeList->edge_data->data, nnz, adj_list); - gdf_column_view(graph->transposedAdjList->offsets, adj_list.rowOffsets, - nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); - gdf_column_view(graph->transposedAdjList->indices, adj_list.colIndices, - nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); - gdf_column_view(graph->transposedAdjList->edge_data, adj_list.edgeWeights, - nullptr, adj_list.nnz, graph->edgeList->edge_data->dtype); - } - else { - - CSR_Result adj_list; - status = ConvertCOOtoCSR((int*)graph->edgeList->dest_indices->data, (int*)graph->edgeList->src_indices->data, nnz, adj_list); - gdf_column_view(graph->transposedAdjList->offsets, adj_list.rowOffsets, - nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); - gdf_column_view(graph->transposedAdjList->indices, adj_list.colIndices, - nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); - } - if (status !=0) { - std::cerr << "Could not generate the adj_list" << std::endl; - return GDF_CUDA_ERROR; - } - } - return GDF_SUCCESS; -} - gdf_error gdf_degree_impl(int n, int e, gdf_column* col_ptr, gdf_column* degree, bool offsets) { if(offsets == true) { dim3 nthreads, nblocks; @@ -475,68 +284,6 @@ gdf_error gdf_pagerank_impl (gdf_graph *graph, return GDF_SUCCESS; } -gdf_error gdf_add_adj_list(gdf_graph *graph) { - if (graph->adjList != nullptr) - return GDF_SUCCESS; - - GDF_REQUIRE( graph->edgeList != nullptr , GDF_INVALID_API_CALL); - GDF_REQUIRE( graph->edgeList->src_indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE ); - - if (graph->edgeList->edge_data != nullptr) { - switch (graph->edgeList->edge_data->dtype) { - case GDF_FLOAT32: return gdf_add_adj_list_impl(graph); - case GDF_FLOAT64: return gdf_add_adj_list_impl(graph); - default: return GDF_UNSUPPORTED_DTYPE; - } - } - else { - return gdf_add_adj_list_impl(graph); - } -} - -gdf_error gdf_add_transposed_adj_list(gdf_graph *graph) { - if (graph->edgeList == nullptr) - gdf_add_edge_list(graph); - - GDF_REQUIRE(graph->edgeList->src_indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(graph->edgeList->dest_indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - - if (graph->edgeList->edge_data != nullptr) { - switch (graph->edgeList->edge_data->dtype) { - case GDF_FLOAT32: return gdf_add_transposed_adj_list_impl(graph); - case GDF_FLOAT64: return gdf_add_transposed_adj_list_impl(graph); - default: return GDF_UNSUPPORTED_DTYPE; - } - } - else { - return gdf_add_transposed_adj_list_impl(graph); - } -} - -gdf_error gdf_delete_adj_list(gdf_graph *graph) { - if (graph->adjList) { - delete graph->adjList; - } - graph->adjList = nullptr; - return GDF_SUCCESS; -} - -gdf_error gdf_delete_edge_list(gdf_graph *graph) { - if (graph->edgeList) { - delete graph->edgeList; - } - graph->edgeList = nullptr; - return GDF_SUCCESS; -} - -gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph) { - if (graph->transposedAdjList) { - delete graph->transposedAdjList; - } - graph->transposedAdjList = nullptr; - return GDF_SUCCESS; -} - gdf_error gdf_pagerank(gdf_graph *graph, gdf_column *pagerank, float alpha, float tolerance, int max_iter, bool has_guess) { // // page rank operates on CSR and can't currently support 64-bit integers. @@ -631,7 +378,7 @@ gdf_error gdf_louvain(gdf_graph *graph, void *final_modularity, void *num_level, template gdf_error gdf_snmg_csrmv_impl (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ - + GDF_REQUIRE( part_offsets != nullptr, GDF_INVALID_API_CALL ); GDF_REQUIRE( off != nullptr, GDF_INVALID_API_CALL ); GDF_REQUIRE( ind != nullptr, GDF_INVALID_API_CALL ); @@ -640,9 +387,9 @@ gdf_error gdf_snmg_csrmv_impl (size_t * part_offsets, gdf_column * off, gdf_colu GDF_REQUIRE( off->size > 0, GDF_INVALID_API_CALL ); GDF_REQUIRE( ind->size > 0, GDF_INVALID_API_CALL ); GDF_REQUIRE( val->size > 0, GDF_INVALID_API_CALL ); - GDF_REQUIRE( ind->size == val->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( off->dtype == ind->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( off->null_count + ind->null_count + val->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( ind->size == val->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( off->dtype == ind->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( off->null_count + ind->null_count + val->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); gdf_error status; auto p = omp_get_num_threads(); @@ -655,9 +402,9 @@ gdf_error gdf_snmg_csrmv_impl (size_t * part_offsets, gdf_column * off, gdf_colu x[i]= static_cast(x_cols[i]->data); } status = cugraph::snmg_csrmv(part_offsets, - static_cast(off->data), - static_cast(ind->data), - static_cast(val->data), + static_cast(off->data), + static_cast(ind->data), + static_cast(val->data), x); return status; } diff --git a/cpp/src/structure/cugraph.cu b/cpp/src/structure/cugraph.cu new file mode 100644 index 00000000000..c5cac9cbc72 --- /dev/null +++ b/cpp/src/structure/cugraph.cu @@ -0,0 +1,276 @@ + /* + * Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved. + * + * NVIDIA CORPORATION and its licensors retain all intellectual property + * and proprietary rights in and to this software, related documentation + * and any modifications thereto. Any use, reproduction, disclosure or + * distribution of this software and related documentation without an express + * license agreement from NVIDIA CORPORATION is strictly prohibited. + * + */ + +// Graph analytics features +// Author: Alex Fender afender@nvidia.com + +#include +#include "graph_utils.cuh" +#include "COOtoCSR.cuh" +#include "utilities/error_utils.h" +#include +#include +#include + +#include + +void cpy_column_view(const gdf_column *in, gdf_column *out) { + if (in != nullptr && out !=nullptr) { + gdf_column_view(out, in->data, in->valid, in->size, in->dtype); + } +} + +gdf_error gdf_adj_list_view(gdf_graph *graph, const gdf_column *offsets, + const gdf_column *indices, const gdf_column *edge_data) { + //This function returns an error if this graph object has at least one graph + //representation to prevent a single object storing two different graphs. + GDF_REQUIRE( ((graph->edgeList == nullptr) && (graph->adjList == nullptr) && + (graph->transposedAdjList == nullptr)), GDF_INVALID_API_CALL); + GDF_REQUIRE( offsets->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( (offsets->dtype == indices->dtype), GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( ((offsets->dtype == GDF_INT32) || (offsets->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( (offsets->size > 0), GDF_DATASET_EMPTY ); + + graph->adjList = new gdf_adj_list; + graph->adjList->offsets = new gdf_column; + graph->adjList->indices = new gdf_column; + graph->adjList->ownership = 0; + + cpy_column_view(offsets, graph->adjList->offsets); + cpy_column_view(indices, graph->adjList->indices); + if (edge_data) { + GDF_REQUIRE( indices->size == edge_data->size, GDF_COLUMN_SIZE_MISMATCH ); + graph->adjList->edge_data = new gdf_column; + cpy_column_view(edge_data, graph->adjList->edge_data); + } + else { + graph->adjList->edge_data = nullptr; + } + return GDF_SUCCESS; +} + +gdf_error gdf_adj_list::get_vertex_identifiers(gdf_column *identifiers) { + GDF_REQUIRE( offsets != nullptr , GDF_INVALID_API_CALL); + GDF_REQUIRE( offsets->data != nullptr , GDF_INVALID_API_CALL); + cugraph::sequence((int)offsets->size-1, (int*)identifiers->data); + return GDF_SUCCESS; +} + +gdf_error gdf_adj_list::get_source_indices (gdf_column *src_indices) { + GDF_REQUIRE( offsets != nullptr , GDF_INVALID_API_CALL); + GDF_REQUIRE( offsets->data != nullptr , GDF_INVALID_API_CALL); + GDF_REQUIRE( src_indices->size == indices->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( src_indices->dtype == indices->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( src_indices->size > 0, GDF_DATASET_EMPTY ); + cugraph::offsets_to_indices((int*)offsets->data, offsets->size-1, (int*)src_indices->data); + + return GDF_SUCCESS; +} + +gdf_error gdf_edge_list_view(gdf_graph *graph, const gdf_column *src_indices, + const gdf_column *dest_indices, const gdf_column *edge_data) { + //This function returns an error if this graph object has at least one graph + //representation to prevent a single object storing two different graphs. + GDF_REQUIRE( ((graph->edgeList == nullptr) && (graph->adjList == nullptr) && + (graph->transposedAdjList == nullptr)), GDF_INVALID_API_CALL); + GDF_REQUIRE( src_indices->size == dest_indices->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( src_indices->dtype == dest_indices->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( ((src_indices->dtype == GDF_INT32) || (src_indices->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( src_indices->size > 0, GDF_DATASET_EMPTY ); + GDF_REQUIRE( src_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( dest_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + + graph->edgeList = new gdf_edge_list; + graph->edgeList->src_indices = new gdf_column; + graph->edgeList->dest_indices = new gdf_column; + graph->edgeList->ownership = 0; + + cpy_column_view(src_indices, graph->edgeList->src_indices); + cpy_column_view(dest_indices, graph->edgeList->dest_indices); + if (edge_data) { + GDF_REQUIRE( src_indices->size == edge_data->size, GDF_COLUMN_SIZE_MISMATCH ); + graph->edgeList->edge_data = new gdf_column; + cpy_column_view(edge_data, graph->edgeList->edge_data); + } + else { + graph->edgeList->edge_data = nullptr; + } + + return GDF_SUCCESS; +} + +template +gdf_error gdf_add_adj_list_impl (gdf_graph *graph) { + if (graph->adjList == nullptr) { + GDF_REQUIRE( graph->edgeList != nullptr , GDF_INVALID_API_CALL); + int nnz = graph->edgeList->src_indices->size, status = 0; + graph->adjList = new gdf_adj_list; + graph->adjList->offsets = new gdf_column; + graph->adjList->indices = new gdf_column; + graph->adjList->ownership = 1; + + if (graph->edgeList->edge_data!= nullptr) { + graph->adjList->edge_data = new gdf_column; + + CSR_Result_Weighted adj_list; + status = ConvertCOOtoCSR_weighted((int*)graph->edgeList->src_indices->data, (int*)graph->edgeList->dest_indices->data, (WT*)graph->edgeList->edge_data->data, nnz, adj_list); + + gdf_column_view(graph->adjList->offsets, adj_list.rowOffsets, + nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); + gdf_column_view(graph->adjList->indices, adj_list.colIndices, + nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); + gdf_column_view(graph->adjList->edge_data, adj_list.edgeWeights, + nullptr, adj_list.nnz, graph->edgeList->edge_data->dtype); + } + else { + CSR_Result adj_list; + status = ConvertCOOtoCSR((int*)graph->edgeList->src_indices->data,(int*)graph->edgeList->dest_indices->data, nnz, adj_list); + gdf_column_view(graph->adjList->offsets, adj_list.rowOffsets, + nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); + gdf_column_view(graph->adjList->indices, adj_list.colIndices, + nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); + } + if (status !=0) { + std::cerr << "Could not generate the adj_list" << std::endl; + return GDF_CUDA_ERROR; + } + } + return GDF_SUCCESS; +} + +gdf_error gdf_add_edge_list (gdf_graph *graph) { + if (graph->edgeList == nullptr) { + GDF_REQUIRE( graph->adjList != nullptr , GDF_INVALID_API_CALL); + int *d_src; + graph->edgeList = new gdf_edge_list; + graph->edgeList->src_indices = new gdf_column; + graph->edgeList->dest_indices = new gdf_column; + graph->edgeList->ownership = 2; + + CUDA_TRY(cudaMallocManaged ((void**)&d_src, sizeof(int) * graph->adjList->indices->size)); + + cugraph::offsets_to_indices((int*)graph->adjList->offsets->data, + graph->adjList->offsets->size-1, + (int*)d_src); + + gdf_column_view(graph->edgeList->src_indices, d_src, + nullptr, graph->adjList->indices->size, graph->adjList->indices->dtype); + cpy_column_view(graph->adjList->indices, graph->edgeList->dest_indices); + + if (graph->adjList->edge_data != nullptr) { + graph->edgeList->edge_data = new gdf_column; + cpy_column_view(graph->adjList->edge_data, graph->edgeList->edge_data); + } + } + return GDF_SUCCESS; +} + + +template +gdf_error gdf_add_transposed_adj_list_impl (gdf_graph *graph) { + if (graph->transposedAdjList == nullptr ) { + GDF_REQUIRE( graph->edgeList != nullptr , GDF_INVALID_API_CALL); + int nnz = graph->edgeList->src_indices->size, status = 0; + graph->transposedAdjList = new gdf_adj_list; + graph->transposedAdjList->offsets = new gdf_column; + graph->transposedAdjList->indices = new gdf_column; + graph->transposedAdjList->ownership = 1; + + if (graph->edgeList->edge_data) { + graph->transposedAdjList->edge_data = new gdf_column; + CSR_Result_Weighted adj_list; + status = ConvertCOOtoCSR_weighted( (int*)graph->edgeList->dest_indices->data, (int*)graph->edgeList->src_indices->data, (WT*)graph->edgeList->edge_data->data, nnz, adj_list); + gdf_column_view(graph->transposedAdjList->offsets, adj_list.rowOffsets, + nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); + gdf_column_view(graph->transposedAdjList->indices, adj_list.colIndices, + nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); + gdf_column_view(graph->transposedAdjList->edge_data, adj_list.edgeWeights, + nullptr, adj_list.nnz, graph->edgeList->edge_data->dtype); + } + else { + + CSR_Result adj_list; + status = ConvertCOOtoCSR((int*)graph->edgeList->dest_indices->data, (int*)graph->edgeList->src_indices->data, nnz, adj_list); + gdf_column_view(graph->transposedAdjList->offsets, adj_list.rowOffsets, + nullptr, adj_list.size+1, graph->edgeList->src_indices->dtype); + gdf_column_view(graph->transposedAdjList->indices, adj_list.colIndices, + nullptr, adj_list.nnz, graph->edgeList->src_indices->dtype); + } + if (status !=0) { + std::cerr << "Could not generate the adj_list" << std::endl; + return GDF_CUDA_ERROR; + } + } + return GDF_SUCCESS; +} + +gdf_error gdf_add_adj_list(gdf_graph *graph) { + if (graph->adjList != nullptr) + return GDF_SUCCESS; + + GDF_REQUIRE( graph->edgeList != nullptr , GDF_INVALID_API_CALL); + GDF_REQUIRE( graph->edgeList->src_indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE ); + + if (graph->edgeList->edge_data != nullptr) { + switch (graph->edgeList->edge_data->dtype) { + case GDF_FLOAT32: return gdf_add_adj_list_impl(graph); + case GDF_FLOAT64: return gdf_add_adj_list_impl(graph); + default: return GDF_UNSUPPORTED_DTYPE; + } + } + else { + return gdf_add_adj_list_impl(graph); + } +} + +gdf_error gdf_add_transposed_adj_list(gdf_graph *graph) { + if (graph->edgeList == nullptr) + gdf_add_edge_list(graph); + + GDF_REQUIRE(graph->edgeList->src_indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(graph->edgeList->dest_indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + + if (graph->edgeList->edge_data != nullptr) { + switch (graph->edgeList->edge_data->dtype) { + case GDF_FLOAT32: return gdf_add_transposed_adj_list_impl(graph); + case GDF_FLOAT64: return gdf_add_transposed_adj_list_impl(graph); + default: return GDF_UNSUPPORTED_DTYPE; + } + } + else { + return gdf_add_transposed_adj_list_impl(graph); + } +} + +gdf_error gdf_delete_adj_list(gdf_graph *graph) { + if (graph->adjList) { + delete graph->adjList; + } + graph->adjList = nullptr; + return GDF_SUCCESS; +} + +gdf_error gdf_delete_edge_list(gdf_graph *graph) { + if (graph->edgeList) { + delete graph->edgeList; + } + graph->edgeList = nullptr; + return GDF_SUCCESS; +} + +gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph) { + if (graph->transposedAdjList) { + delete graph->transposedAdjList; + } + graph->transposedAdjList = nullptr; + return GDF_SUCCESS; +} From 787de1a2b88e7622c1c8217686cf72d379f3351e Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Thu, 9 May 2019 07:26:13 -0700 Subject: [PATCH 03/16] Reorganize source directory --- cpp/CMakeLists.txt | 20 +- cpp/include/nvgraph_gdf.h | 10 +- cpp/src/community/louvain.cu | 65 +++ .../nvgraph_community.cu} | 262 +---------- cpp/src/{ => converters}/COOtoCSR.cuh | 0 cpp/src/converters/nvgraph.cu | 119 +++++ cpp/src/converters/renumber.cu | 116 +++++ cpp/src/{ => converters}/renumber.cuh | 20 +- cpp/src/cugraph.cu | 418 ------------------ cpp/src/{ => link_analysis}/pagerank.cu | 391 +++++++++------- cpp/src/{ => link_prediction}/jaccard.cu | 2 +- cpp/src/{ => link_prediction}/overlap.cu | 4 +- cpp/src/pagerank.cuh | 23 - cpp/src/snmg/blas/snmg_csrmv.cu | 44 ++ cpp/src/snmg/{ => blas}/spmv.cuh | 4 +- cpp/src/snmg/{ => utilities}/snmg_utils.cuh | 0 cpp/src/structure/cugraph.cu | 29 +- cpp/src/tests/renumber/renumber_test.cu | 2 +- cpp/src/{ => traversal}/bfs.cu | 29 +- cpp/src/{ => traversal}/bfs.cuh | 0 cpp/src/{ => traversal}/bfs_kernels.cuh | 0 cpp/src/traversal/nvgraph_sssp.cu | 92 ++++ cpp/src/{ => traversal}/two_hop_neighbors.cu | 0 cpp/src/{ => traversal}/two_hop_neighbors.cuh | 0 cpp/src/utilities/degree.cu | 71 +++ cpp/src/utilities/error_utils.h | 3 +- cpp/src/{ => utilities}/graph_utils.cuh | 0 cpp/src/{ => utilities}/grmat.cu | 0 cpp/src/{ => utilities}/heap.cuh | 0 cpp/src/utilities/nvgraph_error_utils.h | 71 +++ 30 files changed, 906 insertions(+), 889 deletions(-) create mode 100644 cpp/src/community/louvain.cu rename cpp/src/{nvgraph_gdf.cu => community/nvgraph_community.cu} (62%) rename cpp/src/{ => converters}/COOtoCSR.cuh (100%) create mode 100644 cpp/src/converters/nvgraph.cu create mode 100644 cpp/src/converters/renumber.cu rename cpp/src/{ => converters}/renumber.cuh (95%) delete mode 100644 cpp/src/cugraph.cu rename cpp/src/{ => link_analysis}/pagerank.cu (57%) rename cpp/src/{ => link_prediction}/jaccard.cu (99%) rename cpp/src/{ => link_prediction}/overlap.cu (99%) delete mode 100644 cpp/src/pagerank.cuh create mode 100644 cpp/src/snmg/blas/snmg_csrmv.cu rename cpp/src/snmg/{ => blas}/spmv.cuh (93%) rename cpp/src/snmg/{ => utilities}/snmg_utils.cuh (100%) rename cpp/src/{ => traversal}/bfs.cu (92%) rename cpp/src/{ => traversal}/bfs.cuh (100%) rename cpp/src/{ => traversal}/bfs_kernels.cuh (100%) create mode 100644 cpp/src/traversal/nvgraph_sssp.cu rename cpp/src/{ => traversal}/two_hop_neighbors.cu (100%) rename cpp/src/{ => traversal}/two_hop_neighbors.cuh (100%) create mode 100644 cpp/src/utilities/degree.cu rename cpp/src/{ => utilities}/graph_utils.cuh (100%) rename cpp/src/{ => utilities}/grmat.cu (100%) rename cpp/src/{ => utilities}/heap.cuh (100%) create mode 100644 cpp/src/utilities/nvgraph_error_utils.h diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 2de2f5d753a..74115f47e8c 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -229,15 +229,21 @@ link_directories("${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}" # CMAKE_CUDA_IMPLICIT ################################################################################################### # - library targets ------------------------------------------------------------------------------- add_library(cugraph SHARED - src/grmat.cu src/cugraph.cu - src/pagerank.cu - src/bfs.cu - src/jaccard.cu - src/overlap.cu - src/nvgraph_gdf.cu - src/two_hop_neighbors.cu + src/traversal/bfs.cu + src/traversal/nvgraph_sssp.cu + src/traversal/two_hop_neighbors.cu + src/link_analysis/pagerank.cu + src/link_prediction/jaccard.cu + src/link_prediction/overlap.cu src/structure/cugraph.cu + src/snmg/blas/snmg_csrmv.cu + src/community/louvain.cu + src/community/nvgraph_community.cu + src/converters/nvgraph.cu + src/converters/renumber.cu + src/utilities/degree.cu + src/utilities/grmat.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/test_utils.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/error_utils.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/misc_utils.cu diff --git a/cpp/include/nvgraph_gdf.h b/cpp/include/nvgraph_gdf.h index dc262181af6..35e485b44e6 100644 --- a/cpp/include/nvgraph_gdf.h +++ b/cpp/include/nvgraph_gdf.h @@ -21,7 +21,7 @@ #pragma once -//#include +#include #include /** @@ -32,10 +32,10 @@ * @param use_transposed True if we are transposing the input graph while wrapping * @return Error code */ -//gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, -// gdf_graph* gdf_G, -// nvgraphGraphDescr_t * nvgraph_G, -// bool use_transposed = false); +gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, + gdf_graph* gdf_G, + nvgraphGraphDescr_t * nvgraph_G, + bool use_transposed = false); /** * Wrapper function for Nvgraph SSSP algorithm diff --git a/cpp/src/community/louvain.cu b/cpp/src/community/louvain.cu new file mode 100644 index 00000000000..dc77e7efde8 --- /dev/null +++ b/cpp/src/community/louvain.cu @@ -0,0 +1,65 @@ + /* + * Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved. + * + * NVIDIA CORPORATION and its licensors retain all intellectual property + * and proprietary rights in and to this software, related documentation + * and any modifications thereto. Any use, reproduction, disclosure or + * distribution of this software and related documentation without an express + * license agreement from NVIDIA CORPORATION is strictly prohibited. + * + */ + +#include +#include +#include "utilities/error_utils.h" +#include + +template +using Vector = thrust::device_vector>; + +gdf_error gdf_louvain(gdf_graph *graph, void *final_modularity, void *num_level, gdf_column *louvain_parts) { + GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); + gdf_error err = gdf_add_adj_list(graph); + if (err != GDF_SUCCESS) + return err; + + size_t n = graph->adjList->offsets->size - 1; + size_t e = graph->adjList->indices->size; + + void* offsets_ptr = graph->adjList->offsets->data; + void* indices_ptr = graph->adjList->indices->data; + + void* value_ptr; + Vector d_values; + if(graph->adjList->edge_data) { + value_ptr = graph->adjList->edge_data->data; + } + else { + cudaStream_t stream { nullptr }; + rmm_temp_allocator allocator(stream); + d_values.resize(graph->adjList->indices->size); + thrust::fill(thrust::cuda::par(allocator).on(stream), d_values.begin(), d_values.end(), 1.0); + value_ptr = (void * ) thrust::raw_pointer_cast(d_values.data()); + } + + void* louvain_parts_ptr = louvain_parts->data; + + auto gdf_to_cudadtype= [](gdf_column *col){ + cudaDataType_t cuda_dtype; + switch(col->dtype){ + case GDF_INT8: cuda_dtype = CUDA_R_8I; break; + case GDF_INT32: cuda_dtype = CUDA_R_32I; break; + case GDF_FLOAT32: cuda_dtype = CUDA_R_32F; break; + case GDF_FLOAT64: cuda_dtype = CUDA_R_64F; break; + default: throw new std::invalid_argument("Cannot convert data type"); + }return cuda_dtype; + }; + + cudaDataType_t index_type = gdf_to_cudadtype(graph->adjList->indices); + cudaDataType_t val_type = graph->adjList->edge_data? gdf_to_cudadtype(graph->adjList->edge_data): CUDA_R_32F; + + nvgraphLouvain(index_type, val_type, n, e, offsets_ptr, indices_ptr, value_ptr, 1, 0, NULL, + final_modularity, louvain_parts_ptr, num_level); + return GDF_SUCCESS; +} + diff --git a/cpp/src/nvgraph_gdf.cu b/cpp/src/community/nvgraph_community.cu similarity index 62% rename from cpp/src/nvgraph_gdf.cu rename to cpp/src/community/nvgraph_community.cu index 7f493ecafe9..0acdd7a4539 100644 --- a/cpp/src/nvgraph_gdf.cu +++ b/cpp/src/community/nvgraph_community.cu @@ -1,268 +1,23 @@ -// -*-c++-*- - -/* - * Copyright (c) 2019, 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 + /* + * Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved. * - * http://www.apache.org/licenses/LICENSE-2.0 + * NVIDIA CORPORATION and its licensors retain all intellectual property + * and proprietary rights in and to this software, related documentation + * and any modifications thereto. Any use, reproduction, disclosure or + * distribution of this software and related documentation without an express + * license agreement from NVIDIA CORPORATION is strictly prohibited. * - * 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. */ -/** ---------------------------------------------------------------------------* - * @brief Wrapper functions for Nvgraph - * - * @file nvgraph_gdf.cu - * ---------------------------------------------------------------------------**/ +#include #include #include -#include -#include #include "utilities/error_utils.h" - -//RMM: -// - #include template using Vector = thrust::device_vector>; -gdf_error nvgraph2gdf_error(nvgraphStatus_t nvg_stat) { - switch (nvg_stat) { - case NVGRAPH_STATUS_SUCCESS: - return GDF_SUCCESS; - case NVGRAPH_STATUS_NOT_INITIALIZED: - return GDF_INVALID_API_CALL; - case NVGRAPH_STATUS_INVALID_VALUE: - return GDF_INVALID_API_CALL; - case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: - return GDF_UNSUPPORTED_DTYPE; - case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: - return GDF_INVALID_API_CALL; - default: - return GDF_CUDA_ERROR; - } -} - -gdf_error nvgraph2gdf_error_verbose(nvgraphStatus_t nvg_stat) { - switch (nvg_stat) { - case NVGRAPH_STATUS_NOT_INITIALIZED: - std::cerr << "nvGRAPH not initialized"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_ALLOC_FAILED: - std::cerr << "nvGRAPH alloc failed"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_INVALID_VALUE: - std::cerr << "nvGRAPH invalid value"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_ARCH_MISMATCH: - std::cerr << "nvGRAPH arch mismatch"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_MAPPING_ERROR: - std::cerr << "nvGRAPH mapping error"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_EXECUTION_FAILED: - std::cerr << "nvGRAPH execution failed"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_INTERNAL_ERROR: - std::cerr << "nvGRAPH internal error"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: - std::cerr << "nvGRAPH type not supported"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_NOT_CONVERGED: - std::cerr << "nvGRAPH algorithm failed to converge"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: - std::cerr << "nvGRAPH graph type not supported"; - return GDF_CUDA_ERROR; - default: - std::cerr << "Unknown nvGRAPH Status"; - return GDF_CUDA_ERROR; - } -} - -#ifdef VERBOSE -#define NVG_TRY(call) \ -{ \ - if ((call)!=NVGRAPH_STATUS_SUCCESS) \ - return nvgraph2gdf_error_verbose((call)); \ -} -#else -#define NVG_TRY(call) \ -{ \ - nvgraphStatus_t err_code = (call); \ - if (err_code != NVGRAPH_STATUS_SUCCESS) \ - return nvgraph2gdf_error(err_code); \ -} -#endif - -gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, - gdf_graph* gdf_G, - nvgraphGraphDescr_t* nvg_G, - bool use_transposed) { - - // check input - GDF_REQUIRE(!((gdf_G->edgeList == nullptr) && - (gdf_G->adjList == nullptr) && - (gdf_G->transposedAdjList == nullptr)), - GDF_INVALID_API_CALL); - nvgraphTopologyType_t TT; - cudaDataType_t settype; - // create an nvgraph graph handle - NVG_TRY(nvgraphCreateGraphDescr(nvg_handle, nvg_G)); - // setup nvgraph variables - if (use_transposed) { - // convert edgeList to transposedAdjList - if (gdf_G->transposedAdjList == nullptr) { - GDF_TRY(gdf_add_transposed_adj_list(gdf_G)); - } - // using exiting transposedAdjList if it exisits and if adjList is missing - TT = NVGRAPH_CSC_32; - nvgraphCSCTopology32I_st topoData; - topoData.nvertices = gdf_G->transposedAdjList->offsets->size - 1; - topoData.nedges = gdf_G->transposedAdjList->indices->size; - topoData.destination_offsets = (int *) gdf_G->transposedAdjList->offsets->data; - topoData.source_indices = (int *) gdf_G->transposedAdjList->indices->data; - // attach the transposed adj list - NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); - //attach edge values - if (gdf_G->transposedAdjList->edge_data) { - switch (gdf_G->transposedAdjList->edge_data->dtype) { - case GDF_FLOAT32: - settype = CUDA_R_32F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (float * ) gdf_G->transposedAdjList->edge_data->data)) - break; - case GDF_FLOAT64: - settype = CUDA_R_64F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (double * ) gdf_G->transposedAdjList->edge_data->data)) - break; - default: - return GDF_UNSUPPORTED_DTYPE; - } - } - - } - else { - // convert edgeList to adjList - if (gdf_G->adjList == nullptr) { - GDF_TRY(gdf_add_adj_list(gdf_G)); - } - TT = NVGRAPH_CSR_32; - nvgraphCSRTopology32I_st topoData; - topoData.nvertices = gdf_G->adjList->offsets->size - 1; - topoData.nedges = gdf_G->adjList->indices->size; - topoData.source_offsets = (int *) gdf_G->adjList->offsets->data; - topoData.destination_indices = (int *) gdf_G->adjList->indices->data; - - // attach adj list - NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); - //attach edge values - if (gdf_G->adjList->edge_data) { - switch (gdf_G->adjList->edge_data->dtype) { - case GDF_FLOAT32: - settype = CUDA_R_32F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (float * ) gdf_G->adjList->edge_data->data)) - break; - case GDF_FLOAT64: - settype = CUDA_R_64F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (double * ) gdf_G->adjList->edge_data->data)) - break; - default: - return GDF_UNSUPPORTED_DTYPE; - } - } - } - return GDF_SUCCESS; -} - -gdf_error gdf_sssp_nvgraph(gdf_graph *gdf_G, - const int *source_vert, - gdf_column *sssp_distances) { - - GDF_REQUIRE(gdf_G != nullptr, GDF_INVALID_API_CALL); - GDF_REQUIRE(*source_vert >= 0, GDF_INVALID_API_CALL); - GDF_REQUIRE(*source_vert < sssp_distances->size, GDF_INVALID_API_CALL); - GDF_REQUIRE(sssp_distances != nullptr, GDF_INVALID_API_CALL); - GDF_REQUIRE(sssp_distances->data != nullptr, GDF_INVALID_API_CALL); - GDF_REQUIRE(!sssp_distances->valid, GDF_VALIDITY_UNSUPPORTED); - GDF_REQUIRE(sssp_distances->size > 0, GDF_INVALID_API_CALL); - - // init nvgraph - // TODO : time this call - nvgraphHandle_t nvg_handle = 0; - nvgraphGraphDescr_t nvgraph_G = 0; - cudaDataType_t settype; - - NVG_TRY(nvgraphCreate(&nvg_handle)); - GDF_TRY(gdf_createGraph_nvgraph(nvg_handle, gdf_G, &nvgraph_G, true)); - - int sssp_index = 0; - int weight_index = 0; - Vector d_val; - - //RMM: - // - cudaStream_t stream { nullptr }; - rmm_temp_allocator allocator(stream); - if (gdf_G->transposedAdjList->edge_data == nullptr) { - // use a fp32 vector [1,...,1] - settype = CUDA_R_32F; - d_val.resize(gdf_G->transposedAdjList->indices->size); - thrust::fill(thrust::cuda::par(allocator).on(stream), d_val.begin(), d_val.end(), 1.0); - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - nvgraph_G, - weight_index, - settype, - (void * ) thrust::raw_pointer_cast(d_val.data()))); - } - else { - switch (gdf_G->transposedAdjList->edge_data->dtype) { - case GDF_FLOAT32: - settype = CUDA_R_32F; - break; - case GDF_FLOAT64: - settype = CUDA_R_64F; - break; - default: - return GDF_UNSUPPORTED_DTYPE; - } - } - - NVG_TRY(nvgraphAttachVertexData(nvg_handle, nvgraph_G, 0, settype, sssp_distances->data)); - - NVG_TRY(nvgraphSssp(nvg_handle, nvgraph_G, weight_index, source_vert, sssp_index)); - - NVG_TRY(nvgraphDestroyGraphDescr(nvg_handle, nvgraph_G)); - NVG_TRY(nvgraphDestroy(nvg_handle)); - - return GDF_SUCCESS; -} - gdf_error gdf_balancedCutClustering_nvgraph(gdf_graph* gdf_G, const int num_clusters, const int num_eigen_vects, @@ -279,7 +34,6 @@ gdf_error gdf_balancedCutClustering_nvgraph(gdf_graph* gdf_G, // Ensure that the input graph has values GDF_TRY(gdf_add_adj_list(gdf_G)); - //GDF_REQUIRE(gdf_G->adjList->edge_data != nullptr, GDF_INVALID_API_CALL); // Initialize Nvgraph and wrap the graph nvgraphHandle_t nvg_handle = nullptr; diff --git a/cpp/src/COOtoCSR.cuh b/cpp/src/converters/COOtoCSR.cuh similarity index 100% rename from cpp/src/COOtoCSR.cuh rename to cpp/src/converters/COOtoCSR.cuh diff --git a/cpp/src/converters/nvgraph.cu b/cpp/src/converters/nvgraph.cu new file mode 100644 index 00000000000..2eb29566ea4 --- /dev/null +++ b/cpp/src/converters/nvgraph.cu @@ -0,0 +1,119 @@ +/* + * Copyright (c) 2019, 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. + */ +/** ---------------------------------------------------------------------------* + * @brief Wrapper functions for Nvgraph + * + * @file nvgraph_gdf.cu + * ---------------------------------------------------------------------------**/ + +#include +#include +#include "utilities/error_utils.h" + +gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, + gdf_graph* gdf_G, + nvgraphGraphDescr_t* nvg_G, + bool use_transposed) { + + // check input + GDF_REQUIRE(!((gdf_G->edgeList == nullptr) && + (gdf_G->adjList == nullptr) && + (gdf_G->transposedAdjList == nullptr)), + GDF_INVALID_API_CALL); + nvgraphTopologyType_t TT; + cudaDataType_t settype; + // create an nvgraph graph handle + NVG_TRY(nvgraphCreateGraphDescr(nvg_handle, nvg_G)); + // setup nvgraph variables + if (use_transposed) { + // convert edgeList to transposedAdjList + if (gdf_G->transposedAdjList == nullptr) { + GDF_TRY(gdf_add_transposed_adj_list(gdf_G)); + } + // using exiting transposedAdjList if it exisits and if adjList is missing + TT = NVGRAPH_CSC_32; + nvgraphCSCTopology32I_st topoData; + topoData.nvertices = gdf_G->transposedAdjList->offsets->size - 1; + topoData.nedges = gdf_G->transposedAdjList->indices->size; + topoData.destination_offsets = (int *) gdf_G->transposedAdjList->offsets->data; + topoData.source_indices = (int *) gdf_G->transposedAdjList->indices->data; + // attach the transposed adj list + NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); + //attach edge values + if (gdf_G->transposedAdjList->edge_data) { + switch (gdf_G->transposedAdjList->edge_data->dtype) { + case GDF_FLOAT32: + settype = CUDA_R_32F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (float * ) gdf_G->transposedAdjList->edge_data->data)) + break; + case GDF_FLOAT64: + settype = CUDA_R_64F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (double * ) gdf_G->transposedAdjList->edge_data->data)) + break; + default: + return GDF_UNSUPPORTED_DTYPE; + } + } + + } + else { + // convert edgeList to adjList + if (gdf_G->adjList == nullptr) { + GDF_TRY(gdf_add_adj_list(gdf_G)); + } + TT = NVGRAPH_CSR_32; + nvgraphCSRTopology32I_st topoData; + topoData.nvertices = gdf_G->adjList->offsets->size - 1; + topoData.nedges = gdf_G->adjList->indices->size; + topoData.source_offsets = (int *) gdf_G->adjList->offsets->data; + topoData.destination_indices = (int *) gdf_G->adjList->indices->data; + + // attach adj list + NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); + //attach edge values + if (gdf_G->adjList->edge_data) { + switch (gdf_G->adjList->edge_data->dtype) { + case GDF_FLOAT32: + settype = CUDA_R_32F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (float * ) gdf_G->adjList->edge_data->data)) + break; + case GDF_FLOAT64: + settype = CUDA_R_64F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (double * ) gdf_G->adjList->edge_data->data)) + break; + default: + return GDF_UNSUPPORTED_DTYPE; + } + } + } + return GDF_SUCCESS; +} diff --git a/cpp/src/converters/renumber.cu b/cpp/src/converters/renumber.cu new file mode 100644 index 00000000000..d1a7ba87d56 --- /dev/null +++ b/cpp/src/converters/renumber.cu @@ -0,0 +1,116 @@ +/* + * Copyright (c) 2019, 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. + */ + +// Renumber vertices +// Author: Chuck Hastings charlesh@nvidia.com + +#include "renumber.cuh" + +gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, + gdf_column *src_renumbered, gdf_column *dst_renumbered, + gdf_column *numbering_map) { + + GDF_REQUIRE( src->size == dst->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( src->dtype == dst->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( ((src->dtype == GDF_INT32) || (src->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( src->size > 0, GDF_DATASET_EMPTY ); + + // + // TODO: we're currently renumbering without using valid. We need to + // worry about that at some point, but for now we'll just + // copy the valid pointers to the new columns and go from there. + // + cudaStream_t stream{nullptr}; + + size_t src_size = src->size; + size_t new_size; + + // + // TODO: I assume int64_t for output. A few thoughts: + // + // * I could match src->dtype - since if the raw values fit in an int32_t, + // then the renumbered values must fit within an int32_t + // * If new_size < (2^31 - 1) then I could allocate 32-bit integers + // and copy them in order to make the final footprint smaller. + // + // + // NOTE: Forcing match right now - it appears that cugraph is artficially + // forcing the type to be 32 + if (src->dtype == GDF_INT32) { + int32_t *tmp; + + ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(src_renumbered, tmp, src->valid, src->size, src->dtype); + + ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, dst->dtype); + + gdf_error err = cugraph::renumber_vertices(src_size, + (const int32_t *) src->data, + (const int32_t *) dst->data, + (int32_t *) src_renumbered->data, + (int32_t *) dst_renumbered->data, + &new_size, &tmp); + if (err != GDF_SUCCESS) + return err; + + gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); + } else if (src->dtype == GDF_INT64) { + + // + // NOTE: At the moment, we force the renumbered graph to use + // 32-bit integer ids. Since renumbering is going to make + // the vertex range dense, this limits us to 2 billion + // vertices. + // + // The renumbering code supports 64-bit integer generation + // so we can run this with int64_t output if desired... + // but none of the algorithms support that. + // + int64_t *tmp; + ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(src_renumbered, tmp, src->valid, src->size, GDF_INT32); + + ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, GDF_INT32); + + gdf_error err = cugraph::renumber_vertices(src_size, + (const int64_t *) src->data, + (const int64_t *) dst->data, + (int32_t *) src_renumbered->data, + (int32_t *) dst_renumbered->data, + &new_size, &tmp); + if (err != GDF_SUCCESS) + return err; + + // + // If there are too many vertices then the renumbering overflows so we'll + // return an error. + // + if (new_size > 0x7fffffff) { + ALLOC_FREE_TRY(src_renumbered, stream); + ALLOC_FREE_TRY(dst_renumbered, stream); + return GDF_COLUMN_SIZE_TOO_BIG; + } + + gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); + } else { + return GDF_UNSUPPORTED_DTYPE; + } + + return GDF_SUCCESS; +} + diff --git a/cpp/src/renumber.cuh b/cpp/src/converters/renumber.cuh similarity index 95% rename from cpp/src/renumber.cuh rename to cpp/src/converters/renumber.cuh index 0b05135e3cc..a3b4009a60a 100644 --- a/cpp/src/renumber.cuh +++ b/cpp/src/converters/renumber.cuh @@ -34,8 +34,8 @@ #include #include "utilities/error_utils.h" -#include "graph_utils.cuh" -#include "heap.cuh" +#include "utilities/graph_utils.cuh" +#include "utilities/heap.cuh" #include "rmm_utils.h" namespace cugraph { @@ -134,25 +134,27 @@ namespace cugraph { } - __global__ void SetupHash(hash_type hash_size, index_type *hash_bins_start, index_type *hash_bins_end) { + template + __global__ void SetupHash(H hash_size, I *hash_bins_start, I *hash_bins_end) { hash_bins_end[0] = 0; - for (hash_type i = 0 ; i < hash_size ; ++i) { + for (H i = 0 ; i < hash_size ; ++i) { hash_bins_end[i+1] = hash_bins_end[i] + hash_bins_start[i]; } - for (hash_type i = 0 ; i < (hash_size + 1) ; ++i) { + for (H i = 0 ; i < (hash_size + 1) ; ++i) { hash_bins_start[i] = hash_bins_end[i]; } } - __global__ void ComputeBase(hash_type hash_size, index_type *hash_bins_base) { - index_type sum = 0; - for (hash_type i = 0 ; i < hash_size ; ++i) { + template + __global__ void ComputeBase(H hash_size, I *hash_bins_base) { + I sum = 0; + for (H i = 0 ; i < hash_size ; ++i) { sum += hash_bins_base[i]; } hash_bins_base[hash_size] = sum; - for (hash_type i = hash_size ; i > 0 ; --i) { + for (H i = hash_size ; i > 0 ; --i) { hash_bins_base[i-1] = hash_bins_base[i] - hash_bins_base[i-1]; } } diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu deleted file mode 100644 index b277720c96a..00000000000 --- a/cpp/src/cugraph.cu +++ /dev/null @@ -1,418 +0,0 @@ -// -*-c++-*- - - /* - * Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved. - * - * NVIDIA CORPORATION and its licensors retain all intellectual property - * and proprietary rights in and to this software, related documentation - * and any modifications thereto. Any use, reproduction, disclosure or - * distribution of this software and related documentation without an express - * license agreement from NVIDIA CORPORATION is strictly prohibited. - * - */ - -// Graph analytics features -// Author: Alex Fender afender@nvidia.com - -#include -#include "graph_utils.cuh" -#include "pagerank.cuh" -#include "COOtoCSR.cuh" -#include "utilities/error_utils.h" -#include "bfs.cuh" -#include "renumber.cuh" -#include "snmg/spmv.cuh" -#include -#include -#include - -#include - -template -using Vector = thrust::device_vector>; - -/* - * cudf has gdf_column_free and using this is, in general, better design than - * creating our own, but we will keep this as cudf is planning to remove the - * function. cudf plans to redesign cudf::column to fundamentally solve this - * problem, so once they finished the redesign, we need to update this code to - * use their new features. Until that time, we may rely on this as a temporary - * solution. - */ -void gdf_col_delete(gdf_column* col) { - if (col != nullptr) { - auto stream = cudaStream_t{nullptr}; - if (col->data != nullptr) { - ALLOC_FREE_TRY(col->data, stream); - } - if (col->valid != nullptr) { - ALLOC_FREE_TRY(col->valid, stream); - } -#if 0/* Currently, gdf_column_view does not set col_name, and col_name can have - an arbitrary value, so freeing col_name can lead to freeing a ranodom - address. This problem should be cleaned up once cudf finishes - redesigning cudf::column. */ - if (col->col_name != nullptr) { - free(col->col_name); - } -#endif - delete col; - } -} - -void gdf_col_release(gdf_column* col) { - delete col; -} - -gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, - gdf_column *src_renumbered, gdf_column *dst_renumbered, - gdf_column *numbering_map) { - - GDF_REQUIRE( src->size == dst->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( src->dtype == dst->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( ((src->dtype == GDF_INT32) || (src->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( src->size > 0, GDF_DATASET_EMPTY ); - - // - // TODO: we're currently renumbering without using valid. We need to - // worry about that at some point, but for now we'll just - // copy the valid pointers to the new columns and go from there. - // - cudaStream_t stream{nullptr}; - - size_t src_size = src->size; - size_t new_size; - - // - // TODO: I assume int64_t for output. A few thoughts: - // - // * I could match src->dtype - since if the raw values fit in an int32_t, - // then the renumbered values must fit within an int32_t - // * If new_size < (2^31 - 1) then I could allocate 32-bit integers - // and copy them in order to make the final footprint smaller. - // - // - // NOTE: Forcing match right now - it appears that cugraph is artficially - // forcing the type to be 32 - if (src->dtype == GDF_INT32) { - int32_t *tmp; - - ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(src_renumbered, tmp, src->valid, src->size, src->dtype); - - ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, dst->dtype); - - gdf_error err = cugraph::renumber_vertices(src_size, - (const int32_t *) src->data, - (const int32_t *) dst->data, - (int32_t *) src_renumbered->data, - (int32_t *) dst_renumbered->data, - &new_size, &tmp); - if (err != GDF_SUCCESS) - return err; - - gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); - } else if (src->dtype == GDF_INT64) { - - // - // NOTE: At the moment, we force the renumbered graph to use - // 32-bit integer ids. Since renumbering is going to make - // the vertex range dense, this limits us to 2 billion - // vertices. - // - // The renumbering code supports 64-bit integer generation - // so we can run this with int64_t output if desired... - // but none of the algorithms support that. - // - int64_t *tmp; - ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(src_renumbered, tmp, src->valid, src->size, GDF_INT32); - - ALLOC_MANAGED_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, GDF_INT32); - - gdf_error err = cugraph::renumber_vertices(src_size, - (const int64_t *) src->data, - (const int64_t *) dst->data, - (int32_t *) src_renumbered->data, - (int32_t *) dst_renumbered->data, - &new_size, &tmp); - if (err != GDF_SUCCESS) - return err; - - // - // If there are too many vertices then the renumbering overflows so we'll - // return an error. - // - if (new_size > 0x7fffffff) { - ALLOC_FREE_TRY(src_renumbered, stream); - ALLOC_FREE_TRY(dst_renumbered, stream); - return GDF_COLUMN_SIZE_TOO_BIG; - } - - gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); - } else { - return GDF_UNSUPPORTED_DTYPE; - } - - return GDF_SUCCESS; -} - -gdf_error gdf_degree_impl(int n, int e, gdf_column* col_ptr, gdf_column* degree, bool offsets) { - if(offsets == true) { - dim3 nthreads, nblocks; - nthreads.x = min(n, CUDA_MAX_KERNEL_THREADS); - nthreads.y = 1; - nthreads.z = 1; - nblocks.x = min((n + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); - nblocks.y = 1; - nblocks.z = 1; - - switch (col_ptr->dtype) { - case GDF_INT32: cugraph::degree_offsets <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; - default: return GDF_UNSUPPORTED_DTYPE; - } - } - else { - dim3 nthreads, nblocks; - nthreads.x = min(e, CUDA_MAX_KERNEL_THREADS); - nthreads.y = 1; - nthreads.z = 1; - nblocks.x = min((e + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); - nblocks.y = 1; - nblocks.z = 1; - - switch (col_ptr->dtype) { - case GDF_INT32: cugraph::degree_coo <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; - default: return GDF_UNSUPPORTED_DTYPE; - } - } - return GDF_SUCCESS; -} - - -gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x) { - // Calculates the degree of all vertices of the graph - // x = 0: in+out degree - // x = 1: in-degree - // x = 2: out-degree - GDF_REQUIRE(graph->adjList != nullptr || graph->transposedAdjList != nullptr, GDF_INVALID_API_CALL); - int n; - int e; - if(graph->adjList != nullptr) { - n = graph->adjList->offsets->size -1; - e = graph->adjList->indices->size; - } - else { - n = graph->transposedAdjList->offsets->size - 1; - e = graph->transposedAdjList->indices->size; - } - - if(x!=1) { - // Computes out-degree for x=0 and x=2 - if(graph->adjList) - gdf_degree_impl(n, e, graph->adjList->offsets, degree, true); - else - gdf_degree_impl(n, e, graph->transposedAdjList->indices, degree, false); - } - - if(x!=2) { - // Computes in-degree for x=0 and x=1 - if(graph->adjList) - gdf_degree_impl(n, e, graph->adjList->indices, degree, false); - else - gdf_degree_impl(n, e, graph->transposedAdjList->offsets, degree, true); - } - return GDF_SUCCESS; -} - - -template -gdf_error gdf_pagerank_impl (gdf_graph *graph, - gdf_column *pagerank, float alpha = 0.85, - float tolerance = 1e-4, int max_iter = 200, - bool has_guess = false) { - GDF_REQUIRE( graph->edgeList != nullptr, GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( graph->edgeList->src_indices->size == graph->edgeList->dest_indices->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( graph->edgeList->src_indices->dtype == graph->edgeList->dest_indices->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( graph->edgeList->src_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( graph->edgeList->dest_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( pagerank != nullptr , GDF_INVALID_API_CALL ); - GDF_REQUIRE( pagerank->data != nullptr , GDF_INVALID_API_CALL ); - GDF_REQUIRE( pagerank->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( pagerank->size > 0 , GDF_INVALID_API_CALL ); - - int m=pagerank->size, nnz = graph->edgeList->src_indices->size, status = 0; - WT *d_pr, *d_val = nullptr, *d_leaf_vector = nullptr; - WT res = 1.0; - WT *residual = &res; - - if (graph->transposedAdjList == nullptr) { - gdf_add_transposed_adj_list(graph); - } - cudaStream_t stream{nullptr}; - ALLOC_MANAGED_TRY((void**)&d_leaf_vector, sizeof(WT) * m, stream); - ALLOC_MANAGED_TRY((void**)&d_val, sizeof(WT) * nnz , stream); - ALLOC_MANAGED_TRY((void**)&d_pr, sizeof(WT) * m, stream); - - // The templating for HT_matrix_csc_coo assumes that m, nnz and data are all the same type - cugraph::HT_matrix_csc_coo(m, nnz, (int *)graph->transposedAdjList->offsets->data, (int *)graph->transposedAdjList->indices->data, d_val, d_leaf_vector); - - if (has_guess) - { - GDF_REQUIRE( pagerank->data != nullptr, GDF_VALIDITY_UNSUPPORTED ); - cugraph::copy(m, (WT*)pagerank->data, d_pr); - } - - status = cugraph::pagerank( m,nnz, (int*)graph->transposedAdjList->offsets->data, (int*)graph->transposedAdjList->indices->data, - d_val, alpha, d_leaf_vector, false, tolerance, max_iter, d_pr, residual); - - if (status !=0) - switch ( status ) { - case -1: std::cerr<< "Error : bad parameters in Pagerank"<(m, d_pr, (WT*)pagerank->data); - - ALLOC_FREE_TRY(d_val, stream); - ALLOC_FREE_TRY(d_pr, stream); - ALLOC_FREE_TRY(d_leaf_vector, stream); - - return GDF_SUCCESS; -} - -gdf_error gdf_pagerank(gdf_graph *graph, gdf_column *pagerank, float alpha, float tolerance, int max_iter, bool has_guess) { - // - // page rank operates on CSR and can't currently support 64-bit integers. - // - // If csr doesn't exist, create it. Then check type to make sure it is 32-bit. - // - GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); - gdf_error err = gdf_add_adj_list(graph); - if (err != GDF_SUCCESS) - return err; - - GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - - switch (pagerank->dtype) { - case GDF_FLOAT32: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); - case GDF_FLOAT64: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); - default: return GDF_UNSUPPORTED_DTYPE; - } -} - -gdf_error gdf_bfs(gdf_graph *graph, gdf_column *distances, gdf_column *predecessors, int start_vertex, bool directed) { - GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); - gdf_error err = gdf_add_adj_list(graph); - if (err != GDF_SUCCESS) - return err; - GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(distances->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(predecessors->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - - int n = graph->adjList->offsets->size - 1; - int e = graph->adjList->indices->size; - int* offsets_ptr = (int*)graph->adjList->offsets->data; - int* indices_ptr = (int*)graph->adjList->indices->data; - int* distances_ptr = (int*)distances->data; - int* predecessors_ptr = (int*)predecessors->data; - int alpha = 15; - int beta = 18; - - cugraph::Bfs bfs(n, e, offsets_ptr, indices_ptr, directed, alpha, beta); - bfs.configure(distances_ptr, predecessors_ptr, nullptr); - bfs.traverse(start_vertex); - return GDF_SUCCESS; -} - -gdf_error gdf_louvain(gdf_graph *graph, void *final_modularity, void *num_level, gdf_column *louvain_parts) { - GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); - gdf_error err = gdf_add_adj_list(graph); - if (err != GDF_SUCCESS) - return err; - - size_t n = graph->adjList->offsets->size - 1; - size_t e = graph->adjList->indices->size; - - void* offsets_ptr = graph->adjList->offsets->data; - void* indices_ptr = graph->adjList->indices->data; - - void* value_ptr; - Vector d_values; - if(graph->adjList->edge_data) { - value_ptr = graph->adjList->edge_data->data; - } - else { - cudaStream_t stream { nullptr }; - rmm_temp_allocator allocator(stream); - d_values.resize(graph->adjList->indices->size); - thrust::fill(thrust::cuda::par(allocator).on(stream), d_values.begin(), d_values.end(), 1.0); - value_ptr = (void * ) thrust::raw_pointer_cast(d_values.data()); - } - - void* louvain_parts_ptr = louvain_parts->data; - - auto gdf_to_cudadtype= [](gdf_column *col){ - cudaDataType_t cuda_dtype; - switch(col->dtype){ - case GDF_INT8: cuda_dtype = CUDA_R_8I; break; - case GDF_INT32: cuda_dtype = CUDA_R_32I; break; - case GDF_FLOAT32: cuda_dtype = CUDA_R_32F; break; - case GDF_FLOAT64: cuda_dtype = CUDA_R_64F; break; - default: throw new std::invalid_argument("Cannot convert data type"); - }return cuda_dtype; - }; - - cudaDataType_t index_type = gdf_to_cudadtype(graph->adjList->indices); - cudaDataType_t val_type = graph->adjList->edge_data? gdf_to_cudadtype(graph->adjList->edge_data): CUDA_R_32F; - - nvgraphLouvain(index_type, val_type, n, e, offsets_ptr, indices_ptr, value_ptr, 1, 0, NULL, - final_modularity, louvain_parts_ptr, num_level); - return GDF_SUCCESS; -} - -template -gdf_error gdf_snmg_csrmv_impl (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ - - GDF_REQUIRE( part_offsets != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( off != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( ind != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( val != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( x_cols != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( off->size > 0, GDF_INVALID_API_CALL ); - GDF_REQUIRE( ind->size > 0, GDF_INVALID_API_CALL ); - GDF_REQUIRE( val->size > 0, GDF_INVALID_API_CALL ); - GDF_REQUIRE( ind->size == val->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( off->dtype == ind->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( off->null_count + ind->null_count + val->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - - gdf_error status; - auto p = omp_get_num_threads(); - - val_t* x[p]; - for (auto i = 0; i < p; ++i) - { - GDF_REQUIRE( x_cols[i] != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( x_cols[i]->size > 0, GDF_INVALID_API_CALL ); - x[i]= static_cast(x_cols[i]->data); - } - status = cugraph::snmg_csrmv(part_offsets, - static_cast(off->data), - static_cast(ind->data), - static_cast(val->data), - x); - return status; -} - -gdf_error gdf_snmg_csrmv (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ - switch (val->dtype) { - case GDF_FLOAT32: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); - case GDF_FLOAT64: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); - default: return GDF_UNSUPPORTED_DTYPE; - } -} diff --git a/cpp/src/pagerank.cu b/cpp/src/link_analysis/pagerank.cu similarity index 57% rename from cpp/src/pagerank.cu rename to cpp/src/link_analysis/pagerank.cu index 668e19d1bf3..07caf3e69dc 100644 --- a/cpp/src/pagerank.cu +++ b/cpp/src/link_analysis/pagerank.cu @@ -1,163 +1,228 @@ -/* - * Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved. - * - * NVIDIA CORPORATION and its licensors retain all intellectual property - * and proprietary rights in and to this software, related documentation - * and any modifications thereto. Any use, reproduction, disclosure or - * distribution of this software and related documentation without an express - * license agreement from NVIDIA CORPORATION is strictly prohibited. - * - */ - -// Pagerank solver -// Author: Alex Fender afender@nvidia.com - -#include -#include -#include -#include - #include -#include -#include -#include "graph_utils.cuh" -#include "pagerank.cuh" -#include "cub/cub.cuh" -#include -#include - -#include - -namespace cugraph -{ - -#ifdef DEBUG - #define PR_VERBOSE -#endif -template -bool pagerankIteration( IndexType n, IndexType e, IndexType *cscPtr, IndexType *cscInd,ValueType *cscVal, - ValueType alpha, ValueType *a, ValueType *b, float tolerance, int iter, int max_iter, - ValueType * &tmp, void* cub_d_temp_storage, size_t cub_temp_storage_bytes, - ValueType * &pr, ValueType *residual) { - - ValueType dot_res; - cub::DeviceSpmv::CsrMV(cub_d_temp_storage, cub_temp_storage_bytes, cscVal, - cscPtr, cscInd, tmp, pr, - n, n, e); - - scal(n, alpha, pr); - dot_res = dot( n, a, tmp); - axpy(n, dot_res, b, pr); - scal(n, (ValueType)1.0/nrm2(n, pr) , pr); - axpy(n, (ValueType)-1.0, pr, tmp); - *residual = nrm2(n, tmp); - if (*residual < tolerance) - { - scal(n, (ValueType)1.0/nrm1(n,pr), pr); - return true; - } - else - { - if (iter< max_iter) - { - std::swap(pr, tmp); - } - else - { - scal(n, (ValueType)1.0/nrm1(n,pr), pr); - } - return false; - } -} - -template -int pagerank ( IndexType n, IndexType e, IndexType *cscPtr, IndexType *cscInd, ValueType *cscVal, - ValueType alpha, ValueType *a, bool has_guess, float tolerance, int max_iter, - ValueType * &pagerank_vector, ValueType * &residual) { - int max_it, i = 0 ; - float tol; - bool converged = false; - ValueType randomProbability = static_cast( 1.0/n); - ValueType *b=0, *tmp=0; - void* cub_d_temp_storage = NULL; - size_t cub_temp_storage_bytes = 0; - - if (max_iter > 0 ) - max_it = max_iter; - else - max_it = 500; - - if (tolerance == 0.0f) - tol = 1.0E-6f; - else if (tolerance < 1.0f && tolerance > 0.0f) - tol = tolerance; - else - return -1; - - if (alpha <= 0.0f || alpha >= 1.0f) - return -1; - - cudaStream_t stream{nullptr}; - - ALLOC_MANAGED_TRY ((void**)&b, sizeof(ValueType) * n, stream); - ALLOC_MANAGED_TRY ((void**)&tmp, sizeof(ValueType) * n, stream); - cudaCheckError(); - - if (!has_guess) { - fill(n, pagerank_vector, randomProbability); - fill(n, tmp, randomProbability); - } - else { - copy(n, pagerank_vector, tmp); - } - - - fill(n, b, randomProbability); - update_dangling_nodes(n, a, alpha); - - cub::DeviceSpmv::CsrMV(cub_d_temp_storage, cub_temp_storage_bytes, cscVal, - cscPtr, cscInd, tmp, pagerank_vector, n, n, e); - // Allocate temporary storage - ALLOC_MANAGED_TRY ((void**)&cub_d_temp_storage, cub_temp_storage_bytes, stream); - cudaCheckError() - #ifdef PR_VERBOSE - std::stringstream ss; - ss.str(std::string()); - ss <<" ------------------PageRank------------------"<< std::endl; - ss <<" --------------------------------------------"<< std::endl; - ss << std::setw(10) << "Iteration" << std::setw(15) << "Residual" << std::endl; - ss <<" --------------------------------------------"<< std::endl; - std::cout< ( int n, int e, int *cscPtr, int *cscInd,half *cscVal, half alpha, half *a, bool has_guess, float tolerance, int max_iter, half * &pagerank_vector, half * &residual); -template int pagerank ( int n, int e, int *cscPtr, int *cscInd,float *cscVal, float alpha, float *a, bool has_guess, float tolerance, int max_iter, float * &pagerank_vector, float * &residual); -template int pagerank ( int n, int e, int *cscPtr, int *cscInd,double *cscVal, double alpha, double *a, bool has_guess, float tolerance, int max_iter, double * &pagerank_vector, double * &residual); - -} //namespace cugraph + /* + * Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved. + * + * NVIDIA CORPORATION and its licensors retain all intellectual property + * and proprietary rights in and to this software, related documentation + * and any modifications thereto. Any use, reproduction, disclosure or + * distribution of this software and related documentation without an express + * license agreement from NVIDIA CORPORATION is strictly prohibited. + * + */ + +#include +#include "utilities/graph_utils.cuh" +#include "utilities/error_utils.h" +#include "cub/cub.cuh" +#include + +namespace cugraph +{ + +#ifdef DEBUG + #define PR_VERBOSE +#endif +template +bool pagerankIteration( IndexType n, IndexType e, IndexType *cscPtr, IndexType *cscInd,ValueType *cscVal, + ValueType alpha, ValueType *a, ValueType *b, float tolerance, int iter, int max_iter, + ValueType * &tmp, void* cub_d_temp_storage, size_t cub_temp_storage_bytes, + ValueType * &pr, ValueType *residual) { + + ValueType dot_res; + cub::DeviceSpmv::CsrMV(cub_d_temp_storage, cub_temp_storage_bytes, cscVal, + cscPtr, cscInd, tmp, pr, + n, n, e); + + scal(n, alpha, pr); + dot_res = dot( n, a, tmp); + axpy(n, dot_res, b, pr); + scal(n, (ValueType)1.0/nrm2(n, pr) , pr); + axpy(n, (ValueType)-1.0, pr, tmp); + *residual = nrm2(n, tmp); + if (*residual < tolerance) + { + scal(n, (ValueType)1.0/nrm1(n,pr), pr); + return true; + } + else + { + if (iter< max_iter) + { + std::swap(pr, tmp); + } + else + { + scal(n, (ValueType)1.0/nrm1(n,pr), pr); + } + return false; + } +} + +template +int pagerank ( IndexType n, IndexType e, IndexType *cscPtr, IndexType *cscInd, ValueType *cscVal, + ValueType alpha, ValueType *a, bool has_guess, float tolerance, int max_iter, + ValueType * &pagerank_vector, ValueType * &residual) { + int max_it, i = 0 ; + float tol; + bool converged = false; + ValueType randomProbability = static_cast( 1.0/n); + ValueType *b=0, *tmp=0; + void* cub_d_temp_storage = NULL; + size_t cub_temp_storage_bytes = 0; + + if (max_iter > 0 ) + max_it = max_iter; + else + max_it = 500; + + if (tolerance == 0.0f) + tol = 1.0E-6f; + else if (tolerance < 1.0f && tolerance > 0.0f) + tol = tolerance; + else + return -1; + + if (alpha <= 0.0f || alpha >= 1.0f) + return -1; + + cudaStream_t stream{nullptr}; + + ALLOC_MANAGED_TRY ((void**)&b, sizeof(ValueType) * n, stream); + ALLOC_MANAGED_TRY ((void**)&tmp, sizeof(ValueType) * n, stream); + cudaCheckError(); + + if (!has_guess) { + fill(n, pagerank_vector, randomProbability); + fill(n, tmp, randomProbability); + } + else { + copy(n, pagerank_vector, tmp); + } + + + fill(n, b, randomProbability); + update_dangling_nodes(n, a, alpha); + + cub::DeviceSpmv::CsrMV(cub_d_temp_storage, cub_temp_storage_bytes, cscVal, + cscPtr, cscInd, tmp, pagerank_vector, n, n, e); + // Allocate temporary storage + ALLOC_MANAGED_TRY ((void**)&cub_d_temp_storage, cub_temp_storage_bytes, stream); + cudaCheckError() + #ifdef PR_VERBOSE + std::stringstream ss; + ss.str(std::string()); + ss <<" ------------------PageRank------------------"<< std::endl; + ss <<" --------------------------------------------"<< std::endl; + ss << std::setw(10) << "Iteration" << std::setw(15) << "Residual" << std::endl; + ss <<" --------------------------------------------"<< std::endl; + std::cout< ( int n, int e, int *cscPtr, int *cscInd,half *cscVal, half alpha, half *a, bool has_guess, float tolerance, int max_iter, half * &pagerank_vector, half * &residual); +template int pagerank ( int n, int e, int *cscPtr, int *cscInd,float *cscVal, float alpha, float *a, bool has_guess, float tolerance, int max_iter, float * &pagerank_vector, float * &residual); +template int pagerank ( int n, int e, int *cscPtr, int *cscInd,double *cscVal, double alpha, double *a, bool has_guess, float tolerance, int max_iter, double * &pagerank_vector, double * &residual); + +} //namespace cugraph + +template +gdf_error gdf_pagerank_impl (gdf_graph *graph, + gdf_column *pagerank, float alpha = 0.85, + float tolerance = 1e-4, int max_iter = 200, + bool has_guess = false) { + GDF_REQUIRE( graph->edgeList != nullptr, GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( graph->edgeList->src_indices->size == graph->edgeList->dest_indices->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( graph->edgeList->src_indices->dtype == graph->edgeList->dest_indices->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( graph->edgeList->src_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( graph->edgeList->dest_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( pagerank != nullptr , GDF_INVALID_API_CALL ); + GDF_REQUIRE( pagerank->data != nullptr , GDF_INVALID_API_CALL ); + GDF_REQUIRE( pagerank->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( pagerank->size > 0 , GDF_INVALID_API_CALL ); + + int m=pagerank->size, nnz = graph->edgeList->src_indices->size, status = 0; + WT *d_pr, *d_val = nullptr, *d_leaf_vector = nullptr; + WT res = 1.0; + WT *residual = &res; + + if (graph->transposedAdjList == nullptr) { + gdf_add_transposed_adj_list(graph); + } + cudaStream_t stream{nullptr}; + ALLOC_MANAGED_TRY((void**)&d_leaf_vector, sizeof(WT) * m, stream); + ALLOC_MANAGED_TRY((void**)&d_val, sizeof(WT) * nnz , stream); + ALLOC_MANAGED_TRY((void**)&d_pr, sizeof(WT) * m, stream); + + // The templating for HT_matrix_csc_coo assumes that m, nnz and data are all the same type + cugraph::HT_matrix_csc_coo(m, nnz, (int *)graph->transposedAdjList->offsets->data, (int *)graph->transposedAdjList->indices->data, d_val, d_leaf_vector); + + if (has_guess) + { + GDF_REQUIRE( pagerank->data != nullptr, GDF_VALIDITY_UNSUPPORTED ); + cugraph::copy(m, (WT*)pagerank->data, d_pr); + } + + status = cugraph::pagerank( m,nnz, (int*)graph->transposedAdjList->offsets->data, (int*)graph->transposedAdjList->indices->data, + d_val, alpha, d_leaf_vector, false, tolerance, max_iter, d_pr, residual); + + if (status !=0) + switch ( status ) { + case -1: std::cerr<< "Error : bad parameters in Pagerank"<(m, d_pr, (WT*)pagerank->data); + + ALLOC_FREE_TRY(d_val, stream); + ALLOC_FREE_TRY(d_pr, stream); + ALLOC_FREE_TRY(d_leaf_vector, stream); + + return GDF_SUCCESS; +} + +gdf_error gdf_pagerank(gdf_graph *graph, gdf_column *pagerank, float alpha, float tolerance, int max_iter, bool has_guess) { + // + // page rank operates on CSR and can't currently support 64-bit integers. + // + // If csr doesn't exist, create it. Then check type to make sure it is 32-bit. + // + GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); + gdf_error err = gdf_add_adj_list(graph); + if (err != GDF_SUCCESS) + return err; + + GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + + switch (pagerank->dtype) { + case GDF_FLOAT32: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); + case GDF_FLOAT64: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); + default: return GDF_UNSUPPORTED_DTYPE; + } +} diff --git a/cpp/src/jaccard.cu b/cpp/src/link_prediction/jaccard.cu similarity index 99% rename from cpp/src/jaccard.cu rename to cpp/src/link_prediction/jaccard.cu index 91d6206a7e6..06a7e34fcf6 100644 --- a/cpp/src/jaccard.cu +++ b/cpp/src/link_prediction/jaccard.cu @@ -19,7 +19,7 @@ * @file jaccard.cu * ---------------------------------------------------------------------------**/ -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "cugraph.h" #include "rmm_utils.h" #include "utilities/error_utils.h" diff --git a/cpp/src/overlap.cu b/cpp/src/link_prediction/overlap.cu similarity index 99% rename from cpp/src/overlap.cu rename to cpp/src/link_prediction/overlap.cu index 315baf1dac8..8a133104927 100644 --- a/cpp/src/overlap.cu +++ b/cpp/src/link_prediction/overlap.cu @@ -16,10 +16,10 @@ /** ---------------------------------------------------------------------------* * @brief The cugraph Jaccard core functionality * - * @file jaccard.cu + * @file overlap.cu * ---------------------------------------------------------------------------**/ -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "cugraph.h" #include "rmm_utils.h" #include "utilities/error_utils.h" diff --git a/cpp/src/pagerank.cuh b/cpp/src/pagerank.cuh deleted file mode 100644 index d3e1572d3bd..00000000000 --- a/cpp/src/pagerank.cuh +++ /dev/null @@ -1,23 +0,0 @@ -/* - * Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved. - * - * NVIDIA CORPORATION and its licensors retain all intellectual property - * and proprietary rights in and to this software, related documentation - * and any modifications thereto. Any use, reproduction, disclosure or - * distribution of this software and related documentation without an express - * license agreement from NVIDIA CORPORATION is strictly prohibited. - * - */ - -// Pagerank solver -// Author: Alex Fender afender@nvidia.com - -#pragma once -namespace cugraph -{ - -template -int pagerank ( IndexType n, IndexType e, IndexType *cscPtr, IndexType *cscInd,ValueType *cscVal, - ValueType alpha, ValueType *a, bool has_guess, float tolerance, int max_iter, ValueType * &pagerank_vector, ValueType * &residual); - -} //namespace cugraph diff --git a/cpp/src/snmg/blas/snmg_csrmv.cu b/cpp/src/snmg/blas/snmg_csrmv.cu new file mode 100644 index 00000000000..b58135ea2a5 --- /dev/null +++ b/cpp/src/snmg/blas/snmg_csrmv.cu @@ -0,0 +1,44 @@ +#include +#include "utilities/error_utils.h" +#include "snmg/blas/spmv.cuh" + +template +gdf_error gdf_snmg_csrmv_impl (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ + + GDF_REQUIRE( part_offsets != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( off != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( ind != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( val != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( x_cols != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( off->size > 0, GDF_INVALID_API_CALL ); + GDF_REQUIRE( ind->size > 0, GDF_INVALID_API_CALL ); + GDF_REQUIRE( val->size > 0, GDF_INVALID_API_CALL ); + GDF_REQUIRE( ind->size == val->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( off->dtype == ind->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( off->null_count + ind->null_count + val->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + + gdf_error status; + auto p = omp_get_num_threads(); + + val_t* x[p]; + for (auto i = 0; i < p; ++i) + { + GDF_REQUIRE( x_cols[i] != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( x_cols[i]->size > 0, GDF_INVALID_API_CALL ); + x[i]= static_cast(x_cols[i]->data); + } + status = cugraph::snmg_csrmv(part_offsets, + static_cast(off->data), + static_cast(ind->data), + static_cast(val->data), + x); + return status; +} + +gdf_error gdf_snmg_csrmv (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ + switch (val->dtype) { + case GDF_FLOAT32: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); + case GDF_FLOAT64: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); + default: return GDF_UNSUPPORTED_DTYPE; + } +} diff --git a/cpp/src/snmg/spmv.cuh b/cpp/src/snmg/blas/spmv.cuh similarity index 93% rename from cpp/src/snmg/spmv.cuh rename to cpp/src/snmg/blas/spmv.cuh index 3296f0ce1f4..84f7f380426 100644 --- a/cpp/src/snmg/spmv.cuh +++ b/cpp/src/snmg/blas/spmv.cuh @@ -20,8 +20,8 @@ #pragma once #include "cub/cub.cuh" #include -#include "graph_utils.cuh" -#include "snmg_utils.cuh" +#include "utilities/graph_utils.cuh" +#include "snmg/utilities/snmg_utils.cuh" //#define SNMG_DEBUG namespace cugraph diff --git a/cpp/src/snmg/snmg_utils.cuh b/cpp/src/snmg/utilities/snmg_utils.cuh similarity index 100% rename from cpp/src/snmg/snmg_utils.cuh rename to cpp/src/snmg/utilities/snmg_utils.cuh diff --git a/cpp/src/structure/cugraph.cu b/cpp/src/structure/cugraph.cu index c5cac9cbc72..4f379259abc 100644 --- a/cpp/src/structure/cugraph.cu +++ b/cpp/src/structure/cugraph.cu @@ -13,8 +13,8 @@ // Author: Alex Fender afender@nvidia.com #include -#include "graph_utils.cuh" -#include "COOtoCSR.cuh" +#include "utilities/graph_utils.cuh" +#include "converters/COOtoCSR.cuh" #include "utilities/error_utils.h" #include #include @@ -274,3 +274,28 @@ gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph) { graph->transposedAdjList = nullptr; return GDF_SUCCESS; } + +void gdf_col_delete(gdf_column* col) { + if (col != nullptr) { + auto stream = cudaStream_t{nullptr}; + if (col->data != nullptr) { + ALLOC_FREE_TRY(col->data, stream); + } + if (col->valid != nullptr) { + ALLOC_FREE_TRY(col->valid, stream); + } +#if 0/* Currently, gdf_column_view does not set col_name, and col_name can have + an arbitrary value, so freeing col_name can lead to freeing a ranodom + address. This problem should be cleaned up once cudf finishes + redesigning cudf::column. */ + if (col->col_name != nullptr) { + free(col->col_name); + } +#endif + delete col; + } +} + +void gdf_col_release(gdf_column* col) { + delete col; +} diff --git a/cpp/src/tests/renumber/renumber_test.cu b/cpp/src/tests/renumber/renumber_test.cu index cd70e631f3c..436cb311067 100644 --- a/cpp/src/tests/renumber/renumber_test.cu +++ b/cpp/src/tests/renumber/renumber_test.cu @@ -21,7 +21,7 @@ #include "cuda_profiler_api.h" -#include "renumber.cuh" +#include "converters/renumber.cuh" #include "rmm_utils.h" #include diff --git a/cpp/src/bfs.cu b/cpp/src/traversal/bfs.cu similarity index 92% rename from cpp/src/bfs.cu rename to cpp/src/traversal/bfs.cu index 903a514018d..ed6ab871da5 100644 --- a/cpp/src/bfs.cu +++ b/cpp/src/traversal/bfs.cu @@ -9,13 +9,15 @@ * */ +#include #include #include +#include "utilities/error_utils.h" #include "bfs.cuh" #include #include "rmm_utils.h" -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "bfs_kernels.cuh" using namespace bfs_kernels; @@ -469,3 +471,28 @@ namespace cugraph { template class Bfs ; } // end namespace cugraph + +gdf_error gdf_bfs(gdf_graph *graph, gdf_column *distances, gdf_column *predecessors, int start_vertex, bool directed) { + GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); + gdf_error err = gdf_add_adj_list(graph); + if (err != GDF_SUCCESS) + return err; + GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(distances->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(predecessors->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + + int n = graph->adjList->offsets->size - 1; + int e = graph->adjList->indices->size; + int* offsets_ptr = (int*)graph->adjList->offsets->data; + int* indices_ptr = (int*)graph->adjList->indices->data; + int* distances_ptr = (int*)distances->data; + int* predecessors_ptr = (int*)predecessors->data; + int alpha = 15; + int beta = 18; + + cugraph::Bfs bfs(n, e, offsets_ptr, indices_ptr, directed, alpha, beta); + bfs.configure(distances_ptr, predecessors_ptr, nullptr); + bfs.traverse(start_vertex); + return GDF_SUCCESS; +} diff --git a/cpp/src/bfs.cuh b/cpp/src/traversal/bfs.cuh similarity index 100% rename from cpp/src/bfs.cuh rename to cpp/src/traversal/bfs.cuh diff --git a/cpp/src/bfs_kernels.cuh b/cpp/src/traversal/bfs_kernels.cuh similarity index 100% rename from cpp/src/bfs_kernels.cuh rename to cpp/src/traversal/bfs_kernels.cuh diff --git a/cpp/src/traversal/nvgraph_sssp.cu b/cpp/src/traversal/nvgraph_sssp.cu new file mode 100644 index 00000000000..3f4c091faad --- /dev/null +++ b/cpp/src/traversal/nvgraph_sssp.cu @@ -0,0 +1,92 @@ +/* + * Copyright (c) 2019, 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. + */ +/** ---------------------------------------------------------------------------* + * @brief Wrapper functions for Nvgraph + * + * @file nvgraph_sssp.cu + * ---------------------------------------------------------------------------**/ + +#include +#include +#include +#include "utilities/error_utils.h" +#include + +template +using Vector = thrust::device_vector>; + +gdf_error gdf_sssp_nvgraph(gdf_graph *gdf_G, + const int *source_vert, + gdf_column *sssp_distances) { + + GDF_REQUIRE(gdf_G != nullptr, GDF_INVALID_API_CALL); + GDF_REQUIRE(*source_vert >= 0, GDF_INVALID_API_CALL); + GDF_REQUIRE(*source_vert < sssp_distances->size, GDF_INVALID_API_CALL); + GDF_REQUIRE(sssp_distances != nullptr, GDF_INVALID_API_CALL); + GDF_REQUIRE(sssp_distances->data != nullptr, GDF_INVALID_API_CALL); + GDF_REQUIRE(!sssp_distances->valid, GDF_VALIDITY_UNSUPPORTED); + GDF_REQUIRE(sssp_distances->size > 0, GDF_INVALID_API_CALL); + + // init nvgraph + // TODO : time this call + nvgraphHandle_t nvg_handle = 0; + nvgraphGraphDescr_t nvgraph_G = 0; + cudaDataType_t settype; + + NVG_TRY(nvgraphCreate(&nvg_handle)); + GDF_TRY(gdf_createGraph_nvgraph(nvg_handle, gdf_G, &nvgraph_G, true)); + + int sssp_index = 0; + int weight_index = 0; + Vector d_val; + + //RMM: + // + cudaStream_t stream { nullptr }; + rmm_temp_allocator allocator(stream); + if (gdf_G->transposedAdjList->edge_data == nullptr) { + // use a fp32 vector [1,...,1] + settype = CUDA_R_32F; + d_val.resize(gdf_G->transposedAdjList->indices->size); + thrust::fill(thrust::cuda::par(allocator).on(stream), d_val.begin(), d_val.end(), 1.0); + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + nvgraph_G, + weight_index, + settype, + (void * ) thrust::raw_pointer_cast(d_val.data()))); + } + else { + switch (gdf_G->transposedAdjList->edge_data->dtype) { + case GDF_FLOAT32: + settype = CUDA_R_32F; + break; + case GDF_FLOAT64: + settype = CUDA_R_64F; + break; + default: + return GDF_UNSUPPORTED_DTYPE; + } + } + + NVG_TRY(nvgraphAttachVertexData(nvg_handle, nvgraph_G, 0, settype, sssp_distances->data)); + + NVG_TRY(nvgraphSssp(nvg_handle, nvgraph_G, weight_index, source_vert, sssp_index)); + + NVG_TRY(nvgraphDestroyGraphDescr(nvg_handle, nvgraph_G)); + NVG_TRY(nvgraphDestroy(nvg_handle)); + + return GDF_SUCCESS; +} diff --git a/cpp/src/two_hop_neighbors.cu b/cpp/src/traversal/two_hop_neighbors.cu similarity index 100% rename from cpp/src/two_hop_neighbors.cu rename to cpp/src/traversal/two_hop_neighbors.cu diff --git a/cpp/src/two_hop_neighbors.cuh b/cpp/src/traversal/two_hop_neighbors.cuh similarity index 100% rename from cpp/src/two_hop_neighbors.cuh rename to cpp/src/traversal/two_hop_neighbors.cuh diff --git a/cpp/src/utilities/degree.cu b/cpp/src/utilities/degree.cu new file mode 100644 index 00000000000..44b5e210f50 --- /dev/null +++ b/cpp/src/utilities/degree.cu @@ -0,0 +1,71 @@ +#include +#include "utilities/error_utils.h" +#include "utilities/graph_utils.cuh" + +gdf_error gdf_degree_impl(int n, int e, gdf_column* col_ptr, gdf_column* degree, bool offsets) { + if(offsets == true) { + dim3 nthreads, nblocks; + nthreads.x = min(n, CUDA_MAX_KERNEL_THREADS); + nthreads.y = 1; + nthreads.z = 1; + nblocks.x = min((n + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); + nblocks.y = 1; + nblocks.z = 1; + + switch (col_ptr->dtype) { + case GDF_INT32: cugraph::degree_offsets <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; + default: return GDF_UNSUPPORTED_DTYPE; + } + } + else { + dim3 nthreads, nblocks; + nthreads.x = min(e, CUDA_MAX_KERNEL_THREADS); + nthreads.y = 1; + nthreads.z = 1; + nblocks.x = min((e + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); + nblocks.y = 1; + nblocks.z = 1; + + switch (col_ptr->dtype) { + case GDF_INT32: cugraph::degree_coo <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; + default: return GDF_UNSUPPORTED_DTYPE; + } + } + return GDF_SUCCESS; +} + + +gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x) { + // Calculates the degree of all vertices of the graph + // x = 0: in+out degree + // x = 1: in-degree + // x = 2: out-degree + GDF_REQUIRE(graph->adjList != nullptr || graph->transposedAdjList != nullptr, GDF_INVALID_API_CALL); + int n; + int e; + if(graph->adjList != nullptr) { + n = graph->adjList->offsets->size -1; + e = graph->adjList->indices->size; + } + else { + n = graph->transposedAdjList->offsets->size - 1; + e = graph->transposedAdjList->indices->size; + } + + if(x!=1) { + // Computes out-degree for x=0 and x=2 + if(graph->adjList) + gdf_degree_impl(n, e, graph->adjList->offsets, degree, true); + else + gdf_degree_impl(n, e, graph->transposedAdjList->indices, degree, false); + } + + if(x!=2) { + // Computes in-degree for x=0 and x=1 + if(graph->adjList) + gdf_degree_impl(n, e, graph->adjList->indices, degree, false); + else + gdf_degree_impl(n, e, graph->transposedAdjList->offsets, degree, true); + } + return GDF_SUCCESS; +} diff --git a/cpp/src/utilities/error_utils.h b/cpp/src/utilities/error_utils.h index 6b8416da844..06afd1d25f7 100644 --- a/cpp/src/utilities/error_utils.h +++ b/cpp/src/utilities/error_utils.h @@ -25,8 +25,9 @@ #include #include #include +#include "nvgraph_error_utils.h" -#define CUDA_TRY( call ) \ +#define CUDA_TRY( call ) \ { \ cudaError_t cudaStatus = call; \ if ( cudaSuccess != cudaStatus ) \ diff --git a/cpp/src/graph_utils.cuh b/cpp/src/utilities/graph_utils.cuh similarity index 100% rename from cpp/src/graph_utils.cuh rename to cpp/src/utilities/graph_utils.cuh diff --git a/cpp/src/grmat.cu b/cpp/src/utilities/grmat.cu similarity index 100% rename from cpp/src/grmat.cu rename to cpp/src/utilities/grmat.cu diff --git a/cpp/src/heap.cuh b/cpp/src/utilities/heap.cuh similarity index 100% rename from cpp/src/heap.cuh rename to cpp/src/utilities/heap.cuh diff --git a/cpp/src/utilities/nvgraph_error_utils.h b/cpp/src/utilities/nvgraph_error_utils.h new file mode 100644 index 00000000000..8ece5630d43 --- /dev/null +++ b/cpp/src/utilities/nvgraph_error_utils.h @@ -0,0 +1,71 @@ +#ifndef NVGRAPH_ERRORUTILS_H +#define NVGRAPH_ERRORUTILS_H + +#include + +#ifdef VERBOSE +#define NVG_TRY(call) \ +{ \ + nvgraphStatus_t err_code = (call); \ + if (err_code != NVGRAPH_STATUS_SUCCESS) { \ + switch (err_code) { \ + case NVGRAPH_STATUS_SUCCESS: \ + return GDF_SUCCESS; \ + case NVGRAPH_STATUS_NOT_INITIALIZED: \ + return GDF_INVALID_API_CALL; \ + case NVGRAPH_STATUS_INVALID_VALUE: \ + return GDF_INVALID_API_CALL; \ + case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: \ + return GDF_UNSUPPORTED_DTYPE; \ + case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: \ + return GDF_INVALID_API_CALL; \ + default: \ + return GDF_CUDA_ERROR; \ + } \ + } \ +} +#else +#define NVG_TRY(call) \ +{ \ + nvgraphStatus_t err_code = (call); \ + if (err_code != NVGRAPH_STATUS_SUCCESS) { \ + switch (err_code) { \ + case NVGRAPH_STATUS_NOT_INITIALIZED: \ + std::cerr << "nvGRAPH not initialized"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_ALLOC_FAILED: \ + std::cerr << "nvGRAPH alloc failed"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_INVALID_VALUE: \ + std::cerr << "nvGRAPH invalid value"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_ARCH_MISMATCH: \ + std::cerr << "nvGRAPH arch mismatch"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_MAPPING_ERROR: \ + std::cerr << "nvGRAPH mapping error"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_EXECUTION_FAILED: \ + std::cerr << "nvGRAPH execution failed"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_INTERNAL_ERROR: \ + std::cerr << "nvGRAPH internal error"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: \ + std::cerr << "nvGRAPH type not supported"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_NOT_CONVERGED: \ + std::cerr << "nvGRAPH algorithm failed to converge"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: \ + std::cerr << "nvGRAPH graph type not supported"; \ + return GDF_CUDA_ERROR; \ + default: \ + std::cerr << "Unknown nvGRAPH Status"; \ + return GDF_CUDA_ERROR; \ + } \ + } \ +} +#endif + +#endif From 2a783df6d4a25800e362672ab2a4520ba53c6d87 Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Thu, 9 May 2019 09:55:11 -0700 Subject: [PATCH 04/16] Fixed CMakeLists --- cpp/CMakeLists.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 74115f47e8c..057eb32a3c4 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -229,7 +229,6 @@ link_directories("${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}" # CMAKE_CUDA_IMPLICIT ################################################################################################### # - library targets ------------------------------------------------------------------------------- add_library(cugraph SHARED - src/cugraph.cu src/traversal/bfs.cu src/traversal/nvgraph_sssp.cu src/traversal/two_hop_neighbors.cu From 4756eb93975098a21504e784d8e19db7acd56b1e Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Thu, 9 May 2019 10:02:24 -0700 Subject: [PATCH 05/16] Fixed CHANGELOG --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 1246ca02d9f..f677a5bbb6a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -4,6 +4,7 @@ ## Improvements +- PR #286 Reorganized cugraph source directory ## Bug Fixes From 1349c84ad054965daff950ab5358b194679e8552 Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Thu, 9 May 2019 12:11:23 -0700 Subject: [PATCH 06/16] Removed cuda toolkit header dependency nvgraph_gdf no longer include library_types.h --- cpp/include/nvgraph_gdf.h | 14 ------------ cpp/src/community/nvgraph_community.cu | 1 + cpp/src/converters/nvgraph.cu | 1 + cpp/src/converters/nvgraph.cuh | 31 ++++++++++++++++++++++++++ cpp/src/traversal/nvgraph_sssp.cu | 1 + 5 files changed, 34 insertions(+), 14 deletions(-) create mode 100644 cpp/src/converters/nvgraph.cuh diff --git a/cpp/include/nvgraph_gdf.h b/cpp/include/nvgraph_gdf.h index 35e485b44e6..95febeffa26 100644 --- a/cpp/include/nvgraph_gdf.h +++ b/cpp/include/nvgraph_gdf.h @@ -21,22 +21,8 @@ #pragma once -#include #include -/** - * Takes a GDF graph and wraps its data with an Nvgraph graph object. - * @param nvg_handle The Nvgraph handle - * @param gdf_G Pointer to GDF graph object - * @param nvgraph_G Pointer to the Nvgraph graph descriptor - * @param use_transposed True if we are transposing the input graph while wrapping - * @return Error code - */ -gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, - gdf_graph* gdf_G, - nvgraphGraphDescr_t * nvgraph_G, - bool use_transposed = false); - /** * Wrapper function for Nvgraph SSSP algorithm * @param gdf_G Pointer to GDF graph object diff --git a/cpp/src/community/nvgraph_community.cu b/cpp/src/community/nvgraph_community.cu index 0acdd7a4539..8b2ac49dda3 100644 --- a/cpp/src/community/nvgraph_community.cu +++ b/cpp/src/community/nvgraph_community.cu @@ -13,6 +13,7 @@ #include #include #include "utilities/error_utils.h" +#include "converters/nvgraph.cuh" #include template diff --git a/cpp/src/converters/nvgraph.cu b/cpp/src/converters/nvgraph.cu index 2eb29566ea4..cc448b54494 100644 --- a/cpp/src/converters/nvgraph.cu +++ b/cpp/src/converters/nvgraph.cu @@ -22,6 +22,7 @@ #include #include #include "utilities/error_utils.h" +#include "converters/nvgraph.cuh" gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, gdf_graph* gdf_G, diff --git a/cpp/src/converters/nvgraph.cuh b/cpp/src/converters/nvgraph.cuh new file mode 100644 index 00000000000..f7e3388525a --- /dev/null +++ b/cpp/src/converters/nvgraph.cuh @@ -0,0 +1,31 @@ +/* + * Copyright (c) 2019, 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 +#include + +/** + * Takes a GDF graph and wraps its data with an Nvgraph graph object. + * @param nvg_handle The Nvgraph handle + * @param gdf_G Pointer to GDF graph object + * @param nvgraph_G Pointer to the Nvgraph graph descriptor + * @param use_transposed True if we are transposing the input graph while wrapping + * @return Error code + */ +gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, + gdf_graph* gdf_G, + nvgraphGraphDescr_t * nvgraph_G, + bool use_transposed = false); diff --git a/cpp/src/traversal/nvgraph_sssp.cu b/cpp/src/traversal/nvgraph_sssp.cu index 3f4c091faad..73ae4632b45 100644 --- a/cpp/src/traversal/nvgraph_sssp.cu +++ b/cpp/src/traversal/nvgraph_sssp.cu @@ -23,6 +23,7 @@ #include #include #include "utilities/error_utils.h" +#include "converters/nvgraph.cuh" #include template From 7fe7f56218496e0b899530db1e34641419b7aecb Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 18:21:53 -0700 Subject: [PATCH 07/16] Fix file permissions --- cpp/src/bfs.cuh | 0 1 file changed, 0 insertions(+), 0 deletions(-) mode change 100755 => 100644 cpp/src/bfs.cuh diff --git a/cpp/src/bfs.cuh b/cpp/src/bfs.cuh old mode 100755 new mode 100644 From 0db17744b5214f61eadd99c7fd69bd0558f2929a Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 18:39:21 -0700 Subject: [PATCH 08/16] Reorg SNMG --- cpp/CMakeLists.txt | 1 + cpp/src/cugraph.cu | 44 ------------- cpp/src/snmg/blas/spmv.cu | 62 +++++++++++++++++++ cpp/src/snmg/{ => blas}/spmv.cuh | 0 cpp/src/snmg/{ => link_analysis}/pagerank.cuh | 2 +- .../tests/snmg_pagerank/snmg_pagerank_test.cu | 2 +- 6 files changed, 65 insertions(+), 46 deletions(-) create mode 100644 cpp/src/snmg/blas/spmv.cu rename cpp/src/snmg/{ => blas}/spmv.cuh (100%) rename cpp/src/snmg/{ => link_analysis}/pagerank.cuh (96%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 2302085a3e4..53a5e421044 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -237,6 +237,7 @@ add_library(cugraph SHARED src/overlap.cu src/nvgraph_gdf.cu src/two_hop_neighbors.cu + src/snmg/blas/spmv.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/test_utils.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/error_utils.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/misc_utils.cu diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu index 61d1ff2ece1..943a13d3b66 100644 --- a/cpp/src/cugraph.cu +++ b/cpp/src/cugraph.cu @@ -21,8 +21,6 @@ #include "utilities/error_utils.h" #include "bfs.cuh" #include "renumber.cuh" -#include "snmg/spmv.cuh" -#include "snmg/pagerank.cuh" #include #include #include @@ -633,45 +631,3 @@ gdf_error gdf_louvain(gdf_graph *graph, void *final_modularity, void *num_level, final_modularity, louvain_parts_ptr, num_level); return GDF_SUCCESS; } - -template -gdf_error gdf_snmg_csrmv_impl (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ - - GDF_REQUIRE( part_offsets != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( off != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( ind != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( val != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( x_cols != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( off->size > 0, GDF_INVALID_API_CALL ); - GDF_REQUIRE( ind->size > 0, GDF_INVALID_API_CALL ); - GDF_REQUIRE( val->size > 0, GDF_INVALID_API_CALL ); - GDF_REQUIRE( ind->size == val->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( off->dtype == ind->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( off->null_count + ind->null_count + val->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - - auto p = omp_get_num_threads(); - - val_t* x[p]; - for (auto i = 0; i < p; ++i) - { - GDF_REQUIRE( x_cols[i] != nullptr, GDF_INVALID_API_CALL ); - GDF_REQUIRE( x_cols[i]->size > 0, GDF_INVALID_API_CALL ); - x[i]= static_cast(x_cols[i]->data); - } - cugraph::SNMGinfo snmg_env; - cugraph::SNMGcsrmv spmv_solver(snmg_env, part_offsets, - static_cast(off->data), - static_cast(ind->data), - static_cast(val->data), - x); - spmv_solver.run(x); - return GDF_SUCCESS; -} - -gdf_error gdf_snmg_csrmv (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ - switch (val->dtype) { - case GDF_FLOAT32: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); - case GDF_FLOAT64: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); - default: return GDF_UNSUPPORTED_DTYPE; - } -} diff --git a/cpp/src/snmg/blas/spmv.cu b/cpp/src/snmg/blas/spmv.cu new file mode 100644 index 00000000000..c5b369396c7 --- /dev/null +++ b/cpp/src/snmg/blas/spmv.cu @@ -0,0 +1,62 @@ +/* + * Copyright (c) 2019, 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. + */ + +// snmg spmv +// Author: Alex Fender afender@nvidia.com + +#include "spmv.cuh" + +template +gdf_error gdf_snmg_csrmv_impl (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ + + GDF_REQUIRE( part_offsets != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( off != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( ind != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( val != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( x_cols != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( off->size > 0, GDF_INVALID_API_CALL ); + GDF_REQUIRE( ind->size > 0, GDF_INVALID_API_CALL ); + GDF_REQUIRE( val->size > 0, GDF_INVALID_API_CALL ); + GDF_REQUIRE( ind->size == val->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( off->dtype == ind->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( off->null_count + ind->null_count + val->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + + auto p = omp_get_num_threads(); + + val_t* x[p]; + for (auto i = 0; i < p; ++i) + { + GDF_REQUIRE( x_cols[i] != nullptr, GDF_INVALID_API_CALL ); + GDF_REQUIRE( x_cols[i]->size > 0, GDF_INVALID_API_CALL ); + x[i]= static_cast(x_cols[i]->data); + } + cugraph::SNMGinfo snmg_env; + cugraph::SNMGcsrmv spmv_solver(snmg_env, part_offsets, + static_cast(off->data), + static_cast(ind->data), + static_cast(val->data), + x); + spmv_solver.run(x); + return GDF_SUCCESS; +} + +gdf_error gdf_snmg_csrmv (size_t * part_offsets, gdf_column * off, gdf_column * ind, gdf_column * val, gdf_column ** x_cols){ + switch (val->dtype) { + case GDF_FLOAT32: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); + case GDF_FLOAT64: return gdf_snmg_csrmv_impl(part_offsets, off, ind, val, x_cols); + default: return GDF_UNSUPPORTED_DTYPE; + } +} diff --git a/cpp/src/snmg/spmv.cuh b/cpp/src/snmg/blas/spmv.cuh similarity index 100% rename from cpp/src/snmg/spmv.cuh rename to cpp/src/snmg/blas/spmv.cuh diff --git a/cpp/src/snmg/pagerank.cuh b/cpp/src/snmg/link_analysis/pagerank.cuh similarity index 96% rename from cpp/src/snmg/pagerank.cuh rename to cpp/src/snmg/link_analysis/pagerank.cuh index e8c4127ed5b..ea8637a81ac 100644 --- a/cpp/src/snmg/pagerank.cuh +++ b/cpp/src/snmg/link_analysis/pagerank.cuh @@ -22,7 +22,7 @@ #include #include "graph_utils.cuh" #include "snmg/utils.cuh" -#include "snmg/spmv.cuh" +#include "snmg/blas/spmv.cuh" //#define SNMG_DEBUG namespace cugraph diff --git a/cpp/src/tests/snmg_pagerank/snmg_pagerank_test.cu b/cpp/src/tests/snmg_pagerank/snmg_pagerank_test.cu index f4f82c81c48..e65e4267600 100644 --- a/cpp/src/tests/snmg_pagerank/snmg_pagerank_test.cu +++ b/cpp/src/tests/snmg_pagerank/snmg_pagerank_test.cu @@ -21,7 +21,7 @@ #include #include "test_utils.h" #include "snmg_test_utils.h" -#include "snmg/pagerank.cuh" +#include "snmg/link_analysis/pagerank.cuh" //#define SNMG_VERBOSE From bdecc192f524fe38fbebd0c989e08890efb52c08 Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 19:00:07 -0700 Subject: [PATCH 09/16] Reorg Link Analysis --- cpp/CMakeLists.txt | 2 +- cpp/src/cugraph.cu | 86 ------------------------ cpp/src/{ => link_analysis}/pagerank.cu | 87 ++++++++++++++++++++++++- cpp/src/pagerank.cuh | 23 ------- 4 files changed, 87 insertions(+), 111 deletions(-) rename cpp/src/{ => link_analysis}/pagerank.cu (57%) delete mode 100644 cpp/src/pagerank.cuh diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 53a5e421044..32222beb350 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -231,7 +231,7 @@ link_directories("${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}" # CMAKE_CUDA_IMPLICIT add_library(cugraph SHARED src/grmat.cu src/cugraph.cu - src/pagerank.cu + src/link_analysis/pagerank.cu src/bfs.cu src/jaccard.cu src/overlap.cu diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu index 943a13d3b66..4bcb3d09f06 100644 --- a/cpp/src/cugraph.cu +++ b/cpp/src/cugraph.cu @@ -16,7 +16,6 @@ #include #include "graph_utils.cuh" -#include "pagerank.cuh" #include "COOtoCSR.cuh" #include "utilities/error_utils.h" #include "bfs.cuh" @@ -415,70 +414,6 @@ gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x) { } -template -gdf_error gdf_pagerank_impl (gdf_graph *graph, - gdf_column *pagerank, float alpha = 0.85, - float tolerance = 1e-4, int max_iter = 200, - bool has_guess = false) { - GDF_REQUIRE( graph->edgeList != nullptr, GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( graph->edgeList->src_indices->size == graph->edgeList->dest_indices->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( graph->edgeList->src_indices->dtype == graph->edgeList->dest_indices->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( graph->edgeList->src_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( graph->edgeList->dest_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( pagerank != nullptr , GDF_INVALID_API_CALL ); - GDF_REQUIRE( pagerank->data != nullptr , GDF_INVALID_API_CALL ); - GDF_REQUIRE( pagerank->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); - GDF_REQUIRE( pagerank->size > 0 , GDF_INVALID_API_CALL ); - - int m=pagerank->size, nnz = graph->edgeList->src_indices->size, status = 0; - WT *d_pr, *d_val = nullptr, *d_leaf_vector = nullptr; - WT res = 1.0; - WT *residual = &res; - - if (graph->transposedAdjList == nullptr) { - gdf_add_transposed_adj_list(graph); - } - cudaStream_t stream{nullptr}; - ALLOC_TRY((void**)&d_leaf_vector, sizeof(WT) * m, stream); - ALLOC_TRY((void**)&d_val, sizeof(WT) * nnz , stream); -#if 1/* temporary solution till https://github.com/NVlabs/cub/issues/162 is resolved */ - CUDA_TRY(cudaMalloc((void**)&d_pr, sizeof(WT) * m)); -#else - ALLOC_TRY((void**)&d_pr, sizeof(WT) * m, stream); -#endif - - // The templating for HT_matrix_csc_coo assumes that m, nnz and data are all the same type - cugraph::HT_matrix_csc_coo(m, nnz, (int *)graph->transposedAdjList->offsets->data, (int *)graph->transposedAdjList->indices->data, d_val, d_leaf_vector); - - if (has_guess) - { - GDF_REQUIRE( pagerank->data != nullptr, GDF_VALIDITY_UNSUPPORTED ); - cugraph::copy(m, (WT*)pagerank->data, d_pr); - } - - status = cugraph::pagerank( m,nnz, (int*)graph->transposedAdjList->offsets->data, (int*)graph->transposedAdjList->indices->data, - d_val, alpha, d_leaf_vector, false, tolerance, max_iter, d_pr, residual); - - if (status !=0) - switch ( status ) { - case -1: std::cerr<< "Error : bad parameters in Pagerank"<(m, d_pr, (WT*)pagerank->data); - - ALLOC_FREE_TRY(d_val, stream); -#if 1/* temporary solution till https://github.com/NVlabs/cub/issues/162 is resolved */ - CUDA_TRY(cudaFree(d_pr)); -#else - ALLOC_FREE_TRY(d_pr, stream); -#endif - ALLOC_FREE_TRY(d_leaf_vector, stream); - - return GDF_SUCCESS; -} - gdf_error gdf_add_adj_list(gdf_graph *graph) { if (graph->adjList != nullptr) return GDF_SUCCESS; @@ -541,27 +476,6 @@ gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph) { return GDF_SUCCESS; } -gdf_error gdf_pagerank(gdf_graph *graph, gdf_column *pagerank, float alpha, float tolerance, int max_iter, bool has_guess) { - // - // page rank operates on CSR and can't currently support 64-bit integers. - // - // If csr doesn't exist, create it. Then check type to make sure it is 32-bit. - // - GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); - gdf_error err = gdf_add_adj_list(graph); - if (err != GDF_SUCCESS) - return err; - - GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - - switch (pagerank->dtype) { - case GDF_FLOAT32: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); - case GDF_FLOAT64: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); - default: return GDF_UNSUPPORTED_DTYPE; - } -} - gdf_error gdf_bfs(gdf_graph *graph, gdf_column *distances, gdf_column *predecessors, int start_vertex, bool directed) { GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); gdf_error err = gdf_add_adj_list(graph); diff --git a/cpp/src/pagerank.cu b/cpp/src/link_analysis/pagerank.cu similarity index 57% rename from cpp/src/pagerank.cu rename to cpp/src/link_analysis/pagerank.cu index 5d893f320e9..06f24c9b1f2 100644 --- a/cpp/src/pagerank.cu +++ b/cpp/src/link_analysis/pagerank.cu @@ -26,8 +26,8 @@ #include #include "graph_utils.cuh" -#include "pagerank.cuh" #include "utilities/error_utils.h" +#include namespace cugraph { @@ -168,3 +168,88 @@ template int pagerank ( int n, int e, int *cscPtr, int *cscInd,floa template int pagerank ( int n, int e, int *cscPtr, int *cscInd,double *cscVal, double alpha, double *a, bool has_guess, float tolerance, int max_iter, double * &pagerank_vector, double * &residual); } //namespace cugraph + +template +gdf_error gdf_pagerank_impl (gdf_graph *graph, + gdf_column *pagerank, float alpha = 0.85, + float tolerance = 1e-4, int max_iter = 200, + bool has_guess = false) { + GDF_REQUIRE( graph->edgeList != nullptr, GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( graph->edgeList->src_indices->size == graph->edgeList->dest_indices->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( graph->edgeList->src_indices->dtype == graph->edgeList->dest_indices->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( graph->edgeList->src_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( graph->edgeList->dest_indices->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( pagerank != nullptr , GDF_INVALID_API_CALL ); + GDF_REQUIRE( pagerank->data != nullptr , GDF_INVALID_API_CALL ); + GDF_REQUIRE( pagerank->null_count == 0 , GDF_VALIDITY_UNSUPPORTED ); + GDF_REQUIRE( pagerank->size > 0 , GDF_INVALID_API_CALL ); + + int m=pagerank->size, nnz = graph->edgeList->src_indices->size, status = 0; + WT *d_pr, *d_val = nullptr, *d_leaf_vector = nullptr; + WT res = 1.0; + WT *residual = &res; + + if (graph->transposedAdjList == nullptr) { + gdf_add_transposed_adj_list(graph); + } + cudaStream_t stream{nullptr}; + ALLOC_TRY((void**)&d_leaf_vector, sizeof(WT) * m, stream); + ALLOC_TRY((void**)&d_val, sizeof(WT) * nnz , stream); +#if 1/* temporary solution till https://github.com/NVlabs/cub/issues/162 is resolved */ + CUDA_TRY(cudaMalloc((void**)&d_pr, sizeof(WT) * m)); +#else + ALLOC_TRY((void**)&d_pr, sizeof(WT) * m, stream); +#endif + + // The templating for HT_matrix_csc_coo assumes that m, nnz and data are all the same type + cugraph::HT_matrix_csc_coo(m, nnz, (int *)graph->transposedAdjList->offsets->data, (int *)graph->transposedAdjList->indices->data, d_val, d_leaf_vector); + + if (has_guess) + { + GDF_REQUIRE( pagerank->data != nullptr, GDF_VALIDITY_UNSUPPORTED ); + cugraph::copy(m, (WT*)pagerank->data, d_pr); + } + + status = cugraph::pagerank( m,nnz, (int*)graph->transposedAdjList->offsets->data, (int*)graph->transposedAdjList->indices->data, + d_val, alpha, d_leaf_vector, false, tolerance, max_iter, d_pr, residual); + + if (status !=0) + switch ( status ) { + case -1: std::cerr<< "Error : bad parameters in Pagerank"<(m, d_pr, (WT*)pagerank->data); + + ALLOC_FREE_TRY(d_val, stream); +#if 1/* temporary solution till https://github.com/NVlabs/cub/issues/162 is resolved */ + CUDA_TRY(cudaFree(d_pr)); +#else + ALLOC_FREE_TRY(d_pr, stream); +#endif + ALLOC_FREE_TRY(d_leaf_vector, stream); + + return GDF_SUCCESS; +} + +gdf_error gdf_pagerank(gdf_graph *graph, gdf_column *pagerank, float alpha, float tolerance, int max_iter, bool has_guess) { + // + // page rank operates on CSR and can't currently support 64-bit integers. + // + // If csr doesn't exist, create it. Then check type to make sure it is 32-bit. + // + GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); + gdf_error err = gdf_add_adj_list(graph); + if (err != GDF_SUCCESS) + return err; + + GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + + switch (pagerank->dtype) { + case GDF_FLOAT32: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); + case GDF_FLOAT64: return gdf_pagerank_impl(graph, pagerank, alpha, tolerance, max_iter, has_guess); + default: return GDF_UNSUPPORTED_DTYPE; + } +} diff --git a/cpp/src/pagerank.cuh b/cpp/src/pagerank.cuh deleted file mode 100644 index d3e1572d3bd..00000000000 --- a/cpp/src/pagerank.cuh +++ /dev/null @@ -1,23 +0,0 @@ -/* - * Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved. - * - * NVIDIA CORPORATION and its licensors retain all intellectual property - * and proprietary rights in and to this software, related documentation - * and any modifications thereto. Any use, reproduction, disclosure or - * distribution of this software and related documentation without an express - * license agreement from NVIDIA CORPORATION is strictly prohibited. - * - */ - -// Pagerank solver -// Author: Alex Fender afender@nvidia.com - -#pragma once -namespace cugraph -{ - -template -int pagerank ( IndexType n, IndexType e, IndexType *cscPtr, IndexType *cscInd,ValueType *cscVal, - ValueType alpha, ValueType *a, bool has_guess, float tolerance, int max_iter, ValueType * &pagerank_vector, ValueType * &residual); - -} //namespace cugraph From e2573312e0f10e320665eef0bd86758f60df91be Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 19:05:30 -0700 Subject: [PATCH 10/16] Reorg link_prediction --- cpp/CMakeLists.txt | 4 ++-- cpp/src/{ => link_prediction}/jaccard.cu | 0 cpp/src/{ => link_prediction}/overlap.cu | 0 3 files changed, 2 insertions(+), 2 deletions(-) rename cpp/src/{ => link_prediction}/jaccard.cu (100%) rename cpp/src/{ => link_prediction}/overlap.cu (100%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 32222beb350..36a12d4c84e 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -233,8 +233,8 @@ add_library(cugraph SHARED src/cugraph.cu src/link_analysis/pagerank.cu src/bfs.cu - src/jaccard.cu - src/overlap.cu + src/link_prediction/jaccard.cu + src/link_prediction/overlap.cu src/nvgraph_gdf.cu src/two_hop_neighbors.cu src/snmg/blas/spmv.cu diff --git a/cpp/src/jaccard.cu b/cpp/src/link_prediction/jaccard.cu similarity index 100% rename from cpp/src/jaccard.cu rename to cpp/src/link_prediction/jaccard.cu diff --git a/cpp/src/overlap.cu b/cpp/src/link_prediction/overlap.cu similarity index 100% rename from cpp/src/overlap.cu rename to cpp/src/link_prediction/overlap.cu From f81c0052977898d549c09f976742130349dc2301 Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 19:17:06 -0700 Subject: [PATCH 11/16] Reorg traversal --- cpp/CMakeLists.txt | 6 +- cpp/src/converters/nvgraph.cu | 120 ++++++++++ cpp/src/converters/nvgraph.cuh | 31 +++ cpp/src/cugraph.cu | 26 -- cpp/src/nvgraph_gdf.cu | 226 +----------------- cpp/src/{ => traversal}/bfs.cu | 27 +++ cpp/src/{ => traversal}/bfs.cuh | 0 cpp/src/{ => traversal}/bfs_kernels.cuh | 0 cpp/src/traversal/nvgraph_sssp.cu | 88 +++++++ cpp/src/{ => traversal}/two_hop_neighbors.cu | 0 cpp/src/{ => traversal}/two_hop_neighbors.cuh | 0 cpp/src/utilities/error_utils.h | 1 + cpp/src/utilities/nvgraph_error_utils.h | 71 ++++++ 13 files changed, 343 insertions(+), 253 deletions(-) create mode 100644 cpp/src/converters/nvgraph.cu create mode 100644 cpp/src/converters/nvgraph.cuh rename cpp/src/{ => traversal}/bfs.cu (93%) rename cpp/src/{ => traversal}/bfs.cuh (100%) rename cpp/src/{ => traversal}/bfs_kernels.cuh (100%) create mode 100644 cpp/src/traversal/nvgraph_sssp.cu rename cpp/src/{ => traversal}/two_hop_neighbors.cu (100%) rename cpp/src/{ => traversal}/two_hop_neighbors.cuh (100%) create mode 100644 cpp/src/utilities/nvgraph_error_utils.h diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 36a12d4c84e..fa2c27a3024 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -232,11 +232,13 @@ add_library(cugraph SHARED src/grmat.cu src/cugraph.cu src/link_analysis/pagerank.cu - src/bfs.cu + src/traversal/bfs.cu src/link_prediction/jaccard.cu src/link_prediction/overlap.cu + src/converters/nvgraph.cu src/nvgraph_gdf.cu - src/two_hop_neighbors.cu + src/traversal/nvgraph_sssp.cu + src/traversal/two_hop_neighbors.cu src/snmg/blas/spmv.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/test_utils.cu ${CMAKE_CURRENT_BINARY_DIR}/gunrock/gunrock/util/error_utils.cu diff --git a/cpp/src/converters/nvgraph.cu b/cpp/src/converters/nvgraph.cu new file mode 100644 index 00000000000..cc448b54494 --- /dev/null +++ b/cpp/src/converters/nvgraph.cu @@ -0,0 +1,120 @@ +/* + * Copyright (c) 2019, 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. + */ +/** ---------------------------------------------------------------------------* + * @brief Wrapper functions for Nvgraph + * + * @file nvgraph_gdf.cu + * ---------------------------------------------------------------------------**/ + +#include +#include +#include "utilities/error_utils.h" +#include "converters/nvgraph.cuh" + +gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, + gdf_graph* gdf_G, + nvgraphGraphDescr_t* nvg_G, + bool use_transposed) { + + // check input + GDF_REQUIRE(!((gdf_G->edgeList == nullptr) && + (gdf_G->adjList == nullptr) && + (gdf_G->transposedAdjList == nullptr)), + GDF_INVALID_API_CALL); + nvgraphTopologyType_t TT; + cudaDataType_t settype; + // create an nvgraph graph handle + NVG_TRY(nvgraphCreateGraphDescr(nvg_handle, nvg_G)); + // setup nvgraph variables + if (use_transposed) { + // convert edgeList to transposedAdjList + if (gdf_G->transposedAdjList == nullptr) { + GDF_TRY(gdf_add_transposed_adj_list(gdf_G)); + } + // using exiting transposedAdjList if it exisits and if adjList is missing + TT = NVGRAPH_CSC_32; + nvgraphCSCTopology32I_st topoData; + topoData.nvertices = gdf_G->transposedAdjList->offsets->size - 1; + topoData.nedges = gdf_G->transposedAdjList->indices->size; + topoData.destination_offsets = (int *) gdf_G->transposedAdjList->offsets->data; + topoData.source_indices = (int *) gdf_G->transposedAdjList->indices->data; + // attach the transposed adj list + NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); + //attach edge values + if (gdf_G->transposedAdjList->edge_data) { + switch (gdf_G->transposedAdjList->edge_data->dtype) { + case GDF_FLOAT32: + settype = CUDA_R_32F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (float * ) gdf_G->transposedAdjList->edge_data->data)) + break; + case GDF_FLOAT64: + settype = CUDA_R_64F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (double * ) gdf_G->transposedAdjList->edge_data->data)) + break; + default: + return GDF_UNSUPPORTED_DTYPE; + } + } + + } + else { + // convert edgeList to adjList + if (gdf_G->adjList == nullptr) { + GDF_TRY(gdf_add_adj_list(gdf_G)); + } + TT = NVGRAPH_CSR_32; + nvgraphCSRTopology32I_st topoData; + topoData.nvertices = gdf_G->adjList->offsets->size - 1; + topoData.nedges = gdf_G->adjList->indices->size; + topoData.source_offsets = (int *) gdf_G->adjList->offsets->data; + topoData.destination_indices = (int *) gdf_G->adjList->indices->data; + + // attach adj list + NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); + //attach edge values + if (gdf_G->adjList->edge_data) { + switch (gdf_G->adjList->edge_data->dtype) { + case GDF_FLOAT32: + settype = CUDA_R_32F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (float * ) gdf_G->adjList->edge_data->data)) + break; + case GDF_FLOAT64: + settype = CUDA_R_64F; + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + *nvg_G, + 0, + settype, + (double * ) gdf_G->adjList->edge_data->data)) + break; + default: + return GDF_UNSUPPORTED_DTYPE; + } + } + } + return GDF_SUCCESS; +} diff --git a/cpp/src/converters/nvgraph.cuh b/cpp/src/converters/nvgraph.cuh new file mode 100644 index 00000000000..76c1ff97b69 --- /dev/null +++ b/cpp/src/converters/nvgraph.cuh @@ -0,0 +1,31 @@ +/* + * Copyright (c) 2019, 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 +#include + +/** + * Takes a GDF graph and wraps its data with an Nvgraph graph object. + * @param nvg_handle The Nvgraph handle + * @param gdf_G Pointer to GDF graph object + * @param nvgraph_G Pointer to the Nvgraph graph descriptor + * @param use_transposed True if we are transposing the input graph while wrapping + * @return Error code + */ +gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, + gdf_graph* gdf_G, + nvgraphGraphDescr_t * nvgraph_G, +bool use_transposed = false); diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu index 4bcb3d09f06..39911e1b3f6 100644 --- a/cpp/src/cugraph.cu +++ b/cpp/src/cugraph.cu @@ -18,7 +18,6 @@ #include "graph_utils.cuh" #include "COOtoCSR.cuh" #include "utilities/error_utils.h" -#include "bfs.cuh" #include "renumber.cuh" #include #include @@ -476,31 +475,6 @@ gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph) { return GDF_SUCCESS; } -gdf_error gdf_bfs(gdf_graph *graph, gdf_column *distances, gdf_column *predecessors, int start_vertex, bool directed) { - GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); - gdf_error err = gdf_add_adj_list(graph); - if (err != GDF_SUCCESS) - return err; - GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(distances->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - GDF_REQUIRE(predecessors->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); - - int n = graph->adjList->offsets->size - 1; - int e = graph->adjList->indices->size; - int* offsets_ptr = (int*)graph->adjList->offsets->data; - int* indices_ptr = (int*)graph->adjList->indices->data; - int* distances_ptr = (int*)distances->data; - int* predecessors_ptr = (int*)predecessors->data; - int alpha = 15; - int beta = 18; - - cugraph::Bfs bfs(n, e, offsets_ptr, indices_ptr, directed, alpha, beta); - bfs.configure(distances_ptr, predecessors_ptr, nullptr); - bfs.traverse(start_vertex); - return GDF_SUCCESS; -} - gdf_error gdf_louvain(gdf_graph *graph, void *final_modularity, void *num_level, gdf_column *louvain_parts) { GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); gdf_error err = gdf_add_adj_list(graph); diff --git a/cpp/src/nvgraph_gdf.cu b/cpp/src/nvgraph_gdf.cu index 98171511764..9bd5fb468ca 100644 --- a/cpp/src/nvgraph_gdf.cu +++ b/cpp/src/nvgraph_gdf.cu @@ -26,237 +26,13 @@ #include #include #include "utilities/error_utils.h" +#include "converters/nvgraph.cuh" //RMM: // #include -gdf_error nvgraph2gdf_error(nvgraphStatus_t nvg_stat) { - switch (nvg_stat) { - case NVGRAPH_STATUS_SUCCESS: - return GDF_SUCCESS; - case NVGRAPH_STATUS_NOT_INITIALIZED: - return GDF_INVALID_API_CALL; - case NVGRAPH_STATUS_INVALID_VALUE: - return GDF_INVALID_API_CALL; - case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: - return GDF_UNSUPPORTED_DTYPE; - case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: - return GDF_INVALID_API_CALL; - default: - return GDF_CUDA_ERROR; - } -} - -gdf_error nvgraph2gdf_error_verbose(nvgraphStatus_t nvg_stat) { - switch (nvg_stat) { - case NVGRAPH_STATUS_NOT_INITIALIZED: - std::cerr << "nvGRAPH not initialized"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_ALLOC_FAILED: - std::cerr << "nvGRAPH alloc failed"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_INVALID_VALUE: - std::cerr << "nvGRAPH invalid value"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_ARCH_MISMATCH: - std::cerr << "nvGRAPH arch mismatch"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_MAPPING_ERROR: - std::cerr << "nvGRAPH mapping error"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_EXECUTION_FAILED: - std::cerr << "nvGRAPH execution failed"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_INTERNAL_ERROR: - std::cerr << "nvGRAPH internal error"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: - std::cerr << "nvGRAPH type not supported"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_NOT_CONVERGED: - std::cerr << "nvGRAPH algorithm failed to converge"; - return GDF_CUDA_ERROR; - case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: - std::cerr << "nvGRAPH graph type not supported"; - return GDF_CUDA_ERROR; - default: - std::cerr << "Unknown nvGRAPH Status"; - return GDF_CUDA_ERROR; - } -} - -#ifdef VERBOSE -#define NVG_TRY(call) \ -{ \ - if ((call)!=NVGRAPH_STATUS_SUCCESS) \ - return nvgraph2gdf_error_verbose((call)); \ -} -#else -#define NVG_TRY(call) \ -{ \ - nvgraphStatus_t err_code = (call); \ - if (err_code != NVGRAPH_STATUS_SUCCESS) \ - return nvgraph2gdf_error(err_code); \ -} -#endif - -gdf_error gdf_createGraph_nvgraph(nvgraphHandle_t nvg_handle, - gdf_graph* gdf_G, - nvgraphGraphDescr_t* nvg_G, - bool use_transposed) { - - // check input - GDF_REQUIRE(!((gdf_G->edgeList == nullptr) && - (gdf_G->adjList == nullptr) && - (gdf_G->transposedAdjList == nullptr)), - GDF_INVALID_API_CALL); - nvgraphTopologyType_t TT; - cudaDataType_t settype; - // create an nvgraph graph handle - NVG_TRY(nvgraphCreateGraphDescr(nvg_handle, nvg_G)); - // setup nvgraph variables - if (use_transposed) { - // convert edgeList to transposedAdjList - if (gdf_G->transposedAdjList == nullptr) { - GDF_TRY(gdf_add_transposed_adj_list(gdf_G)); - } - // using exiting transposedAdjList if it exisits and if adjList is missing - TT = NVGRAPH_CSC_32; - nvgraphCSCTopology32I_st topoData; - topoData.nvertices = gdf_G->transposedAdjList->offsets->size - 1; - topoData.nedges = gdf_G->transposedAdjList->indices->size; - topoData.destination_offsets = (int *) gdf_G->transposedAdjList->offsets->data; - topoData.source_indices = (int *) gdf_G->transposedAdjList->indices->data; - // attach the transposed adj list - NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); - //attach edge values - if (gdf_G->transposedAdjList->edge_data) { - switch (gdf_G->transposedAdjList->edge_data->dtype) { - case GDF_FLOAT32: - settype = CUDA_R_32F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (float * ) gdf_G->transposedAdjList->edge_data->data)) - break; - case GDF_FLOAT64: - settype = CUDA_R_64F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (double * ) gdf_G->transposedAdjList->edge_data->data)) - break; - default: - return GDF_UNSUPPORTED_DTYPE; - } - } - - } - else { - // convert edgeList to adjList - if (gdf_G->adjList == nullptr) { - GDF_TRY(gdf_add_adj_list(gdf_G)); - } - TT = NVGRAPH_CSR_32; - nvgraphCSRTopology32I_st topoData; - topoData.nvertices = gdf_G->adjList->offsets->size - 1; - topoData.nedges = gdf_G->adjList->indices->size; - topoData.source_offsets = (int *) gdf_G->adjList->offsets->data; - topoData.destination_indices = (int *) gdf_G->adjList->indices->data; - - // attach adj list - NVG_TRY(nvgraphAttachGraphStructure(nvg_handle, *nvg_G, (void * )&topoData, TT)); - //attach edge values - if (gdf_G->adjList->edge_data) { - switch (gdf_G->adjList->edge_data->dtype) { - case GDF_FLOAT32: - settype = CUDA_R_32F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (float * ) gdf_G->adjList->edge_data->data)) - break; - case GDF_FLOAT64: - settype = CUDA_R_64F; - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - *nvg_G, - 0, - settype, - (double * ) gdf_G->adjList->edge_data->data)) - break; - default: - return GDF_UNSUPPORTED_DTYPE; - } - } - } - return GDF_SUCCESS; -} - -gdf_error gdf_sssp_nvgraph(gdf_graph *gdf_G, - const int *source_vert, - gdf_column *sssp_distances) { - GDF_REQUIRE(gdf_G != nullptr, GDF_INVALID_API_CALL); - GDF_REQUIRE(*source_vert >= 0, GDF_INVALID_API_CALL); - GDF_REQUIRE(*source_vert < sssp_distances->size, GDF_INVALID_API_CALL); - GDF_REQUIRE(sssp_distances != nullptr, GDF_INVALID_API_CALL); - GDF_REQUIRE(sssp_distances->data != nullptr, GDF_INVALID_API_CALL); - GDF_REQUIRE(!sssp_distances->valid, GDF_VALIDITY_UNSUPPORTED); - GDF_REQUIRE(sssp_distances->size > 0, GDF_INVALID_API_CALL); - - // init nvgraph - // TODO : time this call - nvgraphHandle_t nvg_handle = 0; - nvgraphGraphDescr_t nvgraph_G = 0; - cudaDataType_t settype; - - NVG_TRY(nvgraphCreate(&nvg_handle)); - GDF_TRY(gdf_createGraph_nvgraph(nvg_handle, gdf_G, &nvgraph_G, true)); - - int sssp_index = 0; - int weight_index = 0; - rmm::device_vector d_val; - - cudaStream_t stream{nullptr}; - - if (gdf_G->transposedAdjList->edge_data == nullptr) { - // use a fp32 vector [1,...,1] - settype = CUDA_R_32F; - d_val.resize(gdf_G->transposedAdjList->indices->size); - thrust::fill(rmm::exec_policy(stream)->on(stream), d_val.begin(), d_val.end(), 1.0); - NVG_TRY(nvgraphAttachEdgeData(nvg_handle, - nvgraph_G, - weight_index, - settype, - (void * ) thrust::raw_pointer_cast(d_val.data()))); - } - else { - switch (gdf_G->transposedAdjList->edge_data->dtype) { - case GDF_FLOAT32: - settype = CUDA_R_32F; - break; - case GDF_FLOAT64: - settype = CUDA_R_64F; - break; - default: - return GDF_UNSUPPORTED_DTYPE; - } - } - - NVG_TRY(nvgraphAttachVertexData(nvg_handle, nvgraph_G, 0, settype, sssp_distances->data)); - - NVG_TRY(nvgraphSssp(nvg_handle, nvgraph_G, weight_index, source_vert, sssp_index)); - - NVG_TRY(nvgraphDestroyGraphDescr(nvg_handle, nvgraph_G)); - NVG_TRY(nvgraphDestroy(nvg_handle)); - - return GDF_SUCCESS; -} - gdf_error gdf_balancedCutClustering_nvgraph(gdf_graph* gdf_G, const int num_clusters, const int num_eigen_vects, diff --git a/cpp/src/bfs.cu b/cpp/src/traversal/bfs.cu similarity index 93% rename from cpp/src/bfs.cu rename to cpp/src/traversal/bfs.cu index 0760b34e8b9..3018b08fcf3 100644 --- a/cpp/src/bfs.cu +++ b/cpp/src/traversal/bfs.cu @@ -9,6 +9,7 @@ * */ +#include #include #include #include "bfs.cuh" @@ -469,3 +470,29 @@ namespace cugraph { template class Bfs ; } // end namespace cugraph + +gdf_error gdf_bfs(gdf_graph *graph, gdf_column *distances, gdf_column *predecessors, int start_vertex, bool directed) { + GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); + gdf_error err = gdf_add_adj_list(graph); + if (err != GDF_SUCCESS) + return err; + GDF_REQUIRE(graph->adjList->offsets->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(graph->adjList->indices->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(distances->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + GDF_REQUIRE(predecessors->dtype == GDF_INT32, GDF_UNSUPPORTED_DTYPE); + + int n = graph->adjList->offsets->size - 1; + int e = graph->adjList->indices->size; + int* offsets_ptr = (int*)graph->adjList->offsets->data; + int* indices_ptr = (int*)graph->adjList->indices->data; + int* distances_ptr = (int*)distances->data; + int* predecessors_ptr = (int*)predecessors->data; + int alpha = 15; + int beta = 18; + + cugraph::Bfs bfs(n, e, offsets_ptr, indices_ptr, directed, alpha, beta); + bfs.configure(distances_ptr, predecessors_ptr, nullptr); + bfs.traverse(start_vertex); + return GDF_SUCCESS; +} + diff --git a/cpp/src/bfs.cuh b/cpp/src/traversal/bfs.cuh similarity index 100% rename from cpp/src/bfs.cuh rename to cpp/src/traversal/bfs.cuh diff --git a/cpp/src/bfs_kernels.cuh b/cpp/src/traversal/bfs_kernels.cuh similarity index 100% rename from cpp/src/bfs_kernels.cuh rename to cpp/src/traversal/bfs_kernels.cuh diff --git a/cpp/src/traversal/nvgraph_sssp.cu b/cpp/src/traversal/nvgraph_sssp.cu new file mode 100644 index 00000000000..fdccfa23c91 --- /dev/null +++ b/cpp/src/traversal/nvgraph_sssp.cu @@ -0,0 +1,88 @@ +/* + * Copyright (c) 2019, 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. + */ + +/** ---------------------------------------------------------------------------* + * @brief Wrapper functions for Nvgraph sssp + * + * @file nvgraph_sssp.cu + * ---------------------------------------------------------------------------**/ + +#include +#include +#include +#include "utilities/error_utils.h" +#include "converters/nvgraph.cuh" +#include + +gdf_error gdf_sssp_nvgraph(gdf_graph *gdf_G, + const int *source_vert, + gdf_column *sssp_distances) { + GDF_REQUIRE(gdf_G != nullptr, GDF_INVALID_API_CALL); + GDF_REQUIRE(*source_vert >= 0, GDF_INVALID_API_CALL); + GDF_REQUIRE(*source_vert < sssp_distances->size, GDF_INVALID_API_CALL); + GDF_REQUIRE(sssp_distances != nullptr, GDF_INVALID_API_CALL); + GDF_REQUIRE(sssp_distances->data != nullptr, GDF_INVALID_API_CALL); + GDF_REQUIRE(!sssp_distances->valid, GDF_VALIDITY_UNSUPPORTED); + GDF_REQUIRE(sssp_distances->size > 0, GDF_INVALID_API_CALL); + + // init nvgraph + // TODO : time this call + nvgraphHandle_t nvg_handle = 0; + nvgraphGraphDescr_t nvgraph_G = 0; + cudaDataType_t settype; + + NVG_TRY(nvgraphCreate(&nvg_handle)); + GDF_TRY(gdf_createGraph_nvgraph(nvg_handle, gdf_G, &nvgraph_G, true)); + + int sssp_index = 0; + int weight_index = 0; + rmm::device_vector d_val; + + cudaStream_t stream{nullptr}; + + if (gdf_G->transposedAdjList->edge_data == nullptr) { + // use a fp32 vector [1,...,1] + settype = CUDA_R_32F; + d_val.resize(gdf_G->transposedAdjList->indices->size); + thrust::fill(rmm::exec_policy(stream)->on(stream), d_val.begin(), d_val.end(), 1.0); + NVG_TRY(nvgraphAttachEdgeData(nvg_handle, + nvgraph_G, + weight_index, + settype, + (void * ) thrust::raw_pointer_cast(d_val.data()))); + } + else { + switch (gdf_G->transposedAdjList->edge_data->dtype) { + case GDF_FLOAT32: + settype = CUDA_R_32F; + break; + case GDF_FLOAT64: + settype = CUDA_R_64F; + break; + default: + return GDF_UNSUPPORTED_DTYPE; + } + } + + NVG_TRY(nvgraphAttachVertexData(nvg_handle, nvgraph_G, 0, settype, sssp_distances->data)); + + NVG_TRY(nvgraphSssp(nvg_handle, nvgraph_G, weight_index, source_vert, sssp_index)); + + NVG_TRY(nvgraphDestroyGraphDescr(nvg_handle, nvgraph_G)); + NVG_TRY(nvgraphDestroy(nvg_handle)); + + return GDF_SUCCESS; +} diff --git a/cpp/src/two_hop_neighbors.cu b/cpp/src/traversal/two_hop_neighbors.cu similarity index 100% rename from cpp/src/two_hop_neighbors.cu rename to cpp/src/traversal/two_hop_neighbors.cu diff --git a/cpp/src/two_hop_neighbors.cuh b/cpp/src/traversal/two_hop_neighbors.cuh similarity index 100% rename from cpp/src/two_hop_neighbors.cuh rename to cpp/src/traversal/two_hop_neighbors.cuh diff --git a/cpp/src/utilities/error_utils.h b/cpp/src/utilities/error_utils.h index 47ca9ef471c..c50feca3a12 100644 --- a/cpp/src/utilities/error_utils.h +++ b/cpp/src/utilities/error_utils.h @@ -28,6 +28,7 @@ #include #include +#include "nvgraph_error_utils.h" #define cudaCheckError() { \ cudaError_t e=cudaGetLastError(); \ diff --git a/cpp/src/utilities/nvgraph_error_utils.h b/cpp/src/utilities/nvgraph_error_utils.h new file mode 100644 index 00000000000..8ece5630d43 --- /dev/null +++ b/cpp/src/utilities/nvgraph_error_utils.h @@ -0,0 +1,71 @@ +#ifndef NVGRAPH_ERRORUTILS_H +#define NVGRAPH_ERRORUTILS_H + +#include + +#ifdef VERBOSE +#define NVG_TRY(call) \ +{ \ + nvgraphStatus_t err_code = (call); \ + if (err_code != NVGRAPH_STATUS_SUCCESS) { \ + switch (err_code) { \ + case NVGRAPH_STATUS_SUCCESS: \ + return GDF_SUCCESS; \ + case NVGRAPH_STATUS_NOT_INITIALIZED: \ + return GDF_INVALID_API_CALL; \ + case NVGRAPH_STATUS_INVALID_VALUE: \ + return GDF_INVALID_API_CALL; \ + case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: \ + return GDF_UNSUPPORTED_DTYPE; \ + case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: \ + return GDF_INVALID_API_CALL; \ + default: \ + return GDF_CUDA_ERROR; \ + } \ + } \ +} +#else +#define NVG_TRY(call) \ +{ \ + nvgraphStatus_t err_code = (call); \ + if (err_code != NVGRAPH_STATUS_SUCCESS) { \ + switch (err_code) { \ + case NVGRAPH_STATUS_NOT_INITIALIZED: \ + std::cerr << "nvGRAPH not initialized"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_ALLOC_FAILED: \ + std::cerr << "nvGRAPH alloc failed"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_INVALID_VALUE: \ + std::cerr << "nvGRAPH invalid value"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_ARCH_MISMATCH: \ + std::cerr << "nvGRAPH arch mismatch"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_MAPPING_ERROR: \ + std::cerr << "nvGRAPH mapping error"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_EXECUTION_FAILED: \ + std::cerr << "nvGRAPH execution failed"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_INTERNAL_ERROR: \ + std::cerr << "nvGRAPH internal error"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_TYPE_NOT_SUPPORTED: \ + std::cerr << "nvGRAPH type not supported"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_NOT_CONVERGED: \ + std::cerr << "nvGRAPH algorithm failed to converge"; \ + return GDF_CUDA_ERROR; \ + case NVGRAPH_STATUS_GRAPH_TYPE_NOT_SUPPORTED: \ + std::cerr << "nvGRAPH graph type not supported"; \ + return GDF_CUDA_ERROR; \ + default: \ + std::cerr << "Unknown nvGRAPH Status"; \ + return GDF_CUDA_ERROR; \ + } \ + } \ +} +#endif + +#endif From 79130f33b88a1f3ef854d12cb3756b6921e7bf2e Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 19:51:30 -0700 Subject: [PATCH 12/16] Reorg converters --- cpp/CMakeLists.txt | 1 + cpp/src/{ => converters}/COOtoCSR.cuh | 0 cpp/src/converters/renumber.cu | 114 ++++++++++++++++++++++++ cpp/src/{ => converters}/renumber.cuh | 16 ++-- cpp/src/cugraph.cu | 98 +------------------- cpp/src/tests/renumber/renumber_test.cu | 2 +- 6 files changed, 127 insertions(+), 104 deletions(-) rename cpp/src/{ => converters}/COOtoCSR.cuh (100%) create mode 100644 cpp/src/converters/renumber.cu rename cpp/src/{ => converters}/renumber.cuh (96%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index fa2c27a3024..04709812c50 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -236,6 +236,7 @@ add_library(cugraph SHARED src/link_prediction/jaccard.cu src/link_prediction/overlap.cu src/converters/nvgraph.cu + src/converters/renumber.cu src/nvgraph_gdf.cu src/traversal/nvgraph_sssp.cu src/traversal/two_hop_neighbors.cu diff --git a/cpp/src/COOtoCSR.cuh b/cpp/src/converters/COOtoCSR.cuh similarity index 100% rename from cpp/src/COOtoCSR.cuh rename to cpp/src/converters/COOtoCSR.cuh diff --git a/cpp/src/converters/renumber.cu b/cpp/src/converters/renumber.cu new file mode 100644 index 00000000000..d7821ab6f55 --- /dev/null +++ b/cpp/src/converters/renumber.cu @@ -0,0 +1,114 @@ +/* + * Copyright (c) 2019, 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. + */ + +// Renumber vertices +// Author: Chuck Hastings charlesh@nvidia.com + +#include "renumber.cuh" + +gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, + gdf_column *src_renumbered, gdf_column *dst_renumbered, + gdf_column *numbering_map) { + GDF_REQUIRE( src->size == dst->size, GDF_COLUMN_SIZE_MISMATCH ); + GDF_REQUIRE( src->dtype == dst->dtype, GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( ((src->dtype == GDF_INT32) || (src->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); + GDF_REQUIRE( src->size > 0, GDF_DATASET_EMPTY ); + + // + // TODO: we're currently renumbering without using valid. We need to + // worry about that at some point, but for now we'll just + // copy the valid pointers to the new columns and go from there. + // + cudaStream_t stream{nullptr}; + + size_t src_size = src->size; + size_t new_size; + + // + // TODO: I assume int64_t for output. A few thoughts: + // + // * I could match src->dtype - since if the raw values fit in an int32_t, + // then the renumbered values must fit within an int32_t + // * If new_size < (2^31 - 1) then I could allocate 32-bit integers + // and copy them in order to make the final footprint smaller. + // + // + // NOTE: Forcing match right now - it appears that cugraph is artficially + // forcing the type to be 32 + if (src->dtype == GDF_INT32) { + int32_t *tmp; + + ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(src_renumbered, tmp, src->valid, src->size, src->dtype); + + ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, dst->dtype); + + gdf_error err = cugraph::renumber_vertices(src_size, + (const int32_t *) src->data, + (const int32_t *) dst->data, + (int32_t *) src_renumbered->data, + (int32_t *) dst_renumbered->data, + &new_size, &tmp); + if (err != GDF_SUCCESS) + return err; + + gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); + } else if (src->dtype == GDF_INT64) { + + // + // NOTE: At the moment, we force the renumbered graph to use + // 32-bit integer ids. Since renumbering is going to make + // the vertex range dense, this limits us to 2 billion + // vertices. + // + // The renumbering code supports 64-bit integer generation + // so we can run this with int64_t output if desired... + // but none of the algorithms support that. + // + int64_t *tmp; + ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(src_renumbered, tmp, src->valid, src->size, GDF_INT32); + + ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); + gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, GDF_INT32); + + gdf_error err = cugraph::renumber_vertices(src_size, + (const int64_t *) src->data, + (const int64_t *) dst->data, + (int32_t *) src_renumbered->data, + (int32_t *) dst_renumbered->data, + &new_size, &tmp); + if (err != GDF_SUCCESS) + return err; + + // + // If there are too many vertices then the renumbering overflows so we'll + // return an error. + // + if (new_size > 0x7fffffff) { + ALLOC_FREE_TRY(src_renumbered, stream); + ALLOC_FREE_TRY(dst_renumbered, stream); + return GDF_COLUMN_SIZE_TOO_BIG; + } + + gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); + } else { + return GDF_UNSUPPORTED_DTYPE; + } + + return GDF_SUCCESS; +} diff --git a/cpp/src/renumber.cuh b/cpp/src/converters/renumber.cuh similarity index 96% rename from cpp/src/renumber.cuh rename to cpp/src/converters/renumber.cuh index b5681491776..d1d0e9f03ba 100644 --- a/cpp/src/renumber.cuh +++ b/cpp/src/converters/renumber.cuh @@ -134,25 +134,27 @@ namespace cugraph { } - __global__ void SetupHash(hash_type hash_size, index_type *hash_bins_start, index_type *hash_bins_end) { + template + __global__ void SetupHash(H hash_size, I *hash_bins_start, I *hash_bins_end) { hash_bins_end[0] = 0; - for (hash_type i = 0 ; i < hash_size ; ++i) { + for (H i = 0 ; i < hash_size ; ++i) { hash_bins_end[i+1] = hash_bins_end[i] + hash_bins_start[i]; } - for (hash_type i = 0 ; i < (hash_size + 1) ; ++i) { + for (H i = 0 ; i < (hash_size + 1) ; ++i) { hash_bins_start[i] = hash_bins_end[i]; } } - __global__ void ComputeBase(hash_type hash_size, index_type *hash_bins_base) { - index_type sum = 0; - for (hash_type i = 0 ; i < hash_size ; ++i) { + template + __global__ void ComputeBase(H hash_size, I *hash_bins_base) { + I sum = 0; + for (H i = 0 ; i < hash_size ; ++i) { sum += hash_bins_base[i]; } hash_bins_base[hash_size] = sum; - for (hash_type i = hash_size ; i > 0 ; --i) { + for (H i = hash_size ; i > 0 ; --i) { hash_bins_base[i-1] = hash_bins_base[i] - hash_bins_base[i-1]; } } diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu index 39911e1b3f6..9e61602e8b7 100644 --- a/cpp/src/cugraph.cu +++ b/cpp/src/cugraph.cu @@ -16,9 +16,9 @@ #include #include "graph_utils.cuh" -#include "COOtoCSR.cuh" +#include "converters/COOtoCSR.cuh" #include "utilities/error_utils.h" -#include "renumber.cuh" +#include "converters/renumber.cuh" #include #include #include @@ -112,100 +112,6 @@ gdf_error gdf_adj_list::get_source_indices (gdf_column *src_indices) { return GDF_SUCCESS; } -gdf_error gdf_renumber_vertices(const gdf_column *src, const gdf_column *dst, - gdf_column *src_renumbered, gdf_column *dst_renumbered, - gdf_column *numbering_map) { - GDF_REQUIRE( src->size == dst->size, GDF_COLUMN_SIZE_MISMATCH ); - GDF_REQUIRE( src->dtype == dst->dtype, GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( ((src->dtype == GDF_INT32) || (src->dtype == GDF_INT64)), GDF_UNSUPPORTED_DTYPE ); - GDF_REQUIRE( src->size > 0, GDF_DATASET_EMPTY ); - - // - // TODO: we're currently renumbering without using valid. We need to - // worry about that at some point, but for now we'll just - // copy the valid pointers to the new columns and go from there. - // - cudaStream_t stream{nullptr}; - - size_t src_size = src->size; - size_t new_size; - - // - // TODO: I assume int64_t for output. A few thoughts: - // - // * I could match src->dtype - since if the raw values fit in an int32_t, - // then the renumbered values must fit within an int32_t - // * If new_size < (2^31 - 1) then I could allocate 32-bit integers - // and copy them in order to make the final footprint smaller. - // - // - // NOTE: Forcing match right now - it appears that cugraph is artficially - // forcing the type to be 32 - if (src->dtype == GDF_INT32) { - int32_t *tmp; - - ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(src_renumbered, tmp, src->valid, src->size, src->dtype); - - ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, dst->dtype); - - gdf_error err = cugraph::renumber_vertices(src_size, - (const int32_t *) src->data, - (const int32_t *) dst->data, - (int32_t *) src_renumbered->data, - (int32_t *) dst_renumbered->data, - &new_size, &tmp); - if (err != GDF_SUCCESS) - return err; - - gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); - } else if (src->dtype == GDF_INT64) { - - // - // NOTE: At the moment, we force the renumbered graph to use - // 32-bit integer ids. Since renumbering is going to make - // the vertex range dense, this limits us to 2 billion - // vertices. - // - // The renumbering code supports 64-bit integer generation - // so we can run this with int64_t output if desired... - // but none of the algorithms support that. - // - int64_t *tmp; - ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(src_renumbered, tmp, src->valid, src->size, GDF_INT32); - - ALLOC_TRY((void**) &tmp, sizeof(int32_t) * src->size, stream); - gdf_column_view(dst_renumbered, tmp, dst->valid, dst->size, GDF_INT32); - - gdf_error err = cugraph::renumber_vertices(src_size, - (const int64_t *) src->data, - (const int64_t *) dst->data, - (int32_t *) src_renumbered->data, - (int32_t *) dst_renumbered->data, - &new_size, &tmp); - if (err != GDF_SUCCESS) - return err; - - // - // If there are too many vertices then the renumbering overflows so we'll - // return an error. - // - if (new_size > 0x7fffffff) { - ALLOC_FREE_TRY(src_renumbered, stream); - ALLOC_FREE_TRY(dst_renumbered, stream); - return GDF_COLUMN_SIZE_TOO_BIG; - } - - gdf_column_view(numbering_map, tmp, nullptr, new_size, src->dtype); - } else { - return GDF_UNSUPPORTED_DTYPE; - } - - return GDF_SUCCESS; -} - gdf_error gdf_edge_list_view(gdf_graph *graph, const gdf_column *src_indices, const gdf_column *dest_indices, const gdf_column *edge_data) { //This function returns an error if this graph object has at least one graph diff --git a/cpp/src/tests/renumber/renumber_test.cu b/cpp/src/tests/renumber/renumber_test.cu index 9a6ed683bce..c982ec71ec1 100644 --- a/cpp/src/tests/renumber/renumber_test.cu +++ b/cpp/src/tests/renumber/renumber_test.cu @@ -21,7 +21,7 @@ #include "cuda_profiler_api.h" -#include "renumber.cuh" +#include "converters/renumber.cuh" #include "rmm_utils.h" #include From 3b0755f191a14cb1648d3902838604fa9e7a0183 Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 19:59:32 -0700 Subject: [PATCH 13/16] Reorg community --- cpp/CMakeLists.txt | 2 +- cpp/src/{ => community}/nvgraph_gdf.cu | 46 ++++++++++++++++++++++++++ cpp/src/cugraph.cu | 45 ------------------------- 3 files changed, 47 insertions(+), 46 deletions(-) rename cpp/src/{ => community}/nvgraph_gdf.cu (89%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 04709812c50..f357af04f9c 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -237,7 +237,7 @@ add_library(cugraph SHARED src/link_prediction/overlap.cu src/converters/nvgraph.cu src/converters/renumber.cu - src/nvgraph_gdf.cu + src/community/nvgraph_gdf.cu src/traversal/nvgraph_sssp.cu src/traversal/two_hop_neighbors.cu src/snmg/blas/spmv.cu diff --git a/cpp/src/nvgraph_gdf.cu b/cpp/src/community/nvgraph_gdf.cu similarity index 89% rename from cpp/src/nvgraph_gdf.cu rename to cpp/src/community/nvgraph_gdf.cu index 9bd5fb468ca..4e605fb91f3 100644 --- a/cpp/src/nvgraph_gdf.cu +++ b/cpp/src/community/nvgraph_gdf.cu @@ -21,6 +21,7 @@ * @file nvgraph_gdf.cu * ---------------------------------------------------------------------------**/ +#include #include #include #include @@ -368,3 +369,48 @@ gdf_error gdf_triangle_count_nvgraph(gdf_graph* G, uint64_t* result) { NVG_TRY(nvgraphTriangleCount(nvg_handle, nvg_G, result)); return GDF_SUCCESS; } + +gdf_error gdf_louvain(gdf_graph *graph, void *final_modularity, void *num_level, gdf_column *louvain_parts) { + GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); + gdf_error err = gdf_add_adj_list(graph); + if (err != GDF_SUCCESS) + return err; + + size_t n = graph->adjList->offsets->size - 1; + size_t e = graph->adjList->indices->size; + + void* offsets_ptr = graph->adjList->offsets->data; + void* indices_ptr = graph->adjList->indices->data; + + void* value_ptr; + rmm::device_vector d_values; + if(graph->adjList->edge_data) { + value_ptr = graph->adjList->edge_data->data; + } + else { + cudaStream_t stream {nullptr}; + d_values.resize(graph->adjList->indices->size); + thrust::fill(rmm::exec_policy(stream)->on(stream), d_values.begin(), d_values.end(), 1.0); + value_ptr = (void * ) thrust::raw_pointer_cast(d_values.data()); + } + + void* louvain_parts_ptr = louvain_parts->data; + + auto gdf_to_cudadtype= [](gdf_column *col){ + cudaDataType_t cuda_dtype; + switch(col->dtype){ + case GDF_INT8: cuda_dtype = CUDA_R_8I; break; + case GDF_INT32: cuda_dtype = CUDA_R_32I; break; + case GDF_FLOAT32: cuda_dtype = CUDA_R_32F; break; + case GDF_FLOAT64: cuda_dtype = CUDA_R_64F; break; + default: throw new std::invalid_argument("Cannot convert data type"); + }return cuda_dtype; + }; + + cudaDataType_t index_type = gdf_to_cudadtype(graph->adjList->indices); + cudaDataType_t val_type = graph->adjList->edge_data? gdf_to_cudadtype(graph->adjList->edge_data): CUDA_R_32F; + + nvgraphLouvain(index_type, val_type, n, e, offsets_ptr, indices_ptr, value_ptr, 1, 0, NULL, + final_modularity, louvain_parts_ptr, num_level); + return GDF_SUCCESS; +} diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu index 9e61602e8b7..505f84ed8d4 100644 --- a/cpp/src/cugraph.cu +++ b/cpp/src/cugraph.cu @@ -380,48 +380,3 @@ gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph) { graph->transposedAdjList = nullptr; return GDF_SUCCESS; } - -gdf_error gdf_louvain(gdf_graph *graph, void *final_modularity, void *num_level, gdf_column *louvain_parts) { - GDF_REQUIRE(graph->adjList != nullptr || graph->edgeList != nullptr, GDF_INVALID_API_CALL); - gdf_error err = gdf_add_adj_list(graph); - if (err != GDF_SUCCESS) - return err; - - size_t n = graph->adjList->offsets->size - 1; - size_t e = graph->adjList->indices->size; - - void* offsets_ptr = graph->adjList->offsets->data; - void* indices_ptr = graph->adjList->indices->data; - - void* value_ptr; - rmm::device_vector d_values; - if(graph->adjList->edge_data) { - value_ptr = graph->adjList->edge_data->data; - } - else { - cudaStream_t stream {nullptr}; - d_values.resize(graph->adjList->indices->size); - thrust::fill(rmm::exec_policy(stream)->on(stream), d_values.begin(), d_values.end(), 1.0); - value_ptr = (void * ) thrust::raw_pointer_cast(d_values.data()); - } - - void* louvain_parts_ptr = louvain_parts->data; - - auto gdf_to_cudadtype= [](gdf_column *col){ - cudaDataType_t cuda_dtype; - switch(col->dtype){ - case GDF_INT8: cuda_dtype = CUDA_R_8I; break; - case GDF_INT32: cuda_dtype = CUDA_R_32I; break; - case GDF_FLOAT32: cuda_dtype = CUDA_R_32F; break; - case GDF_FLOAT64: cuda_dtype = CUDA_R_64F; break; - default: throw new std::invalid_argument("Cannot convert data type"); - }return cuda_dtype; - }; - - cudaDataType_t index_type = gdf_to_cudadtype(graph->adjList->indices); - cudaDataType_t val_type = graph->adjList->edge_data? gdf_to_cudadtype(graph->adjList->edge_data): CUDA_R_32F; - - nvgraphLouvain(index_type, val_type, n, e, offsets_ptr, indices_ptr, value_ptr, 1, 0, NULL, - final_modularity, louvain_parts_ptr, num_level); - return GDF_SUCCESS; -} From 7bdb0dadce52c71c319d3366a140f3d8a2bfe44a Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 21:15:21 -0700 Subject: [PATCH 14/16] Reorg utilities --- cpp/CMakeLists.txt | 3 +- cpp/src/converters/renumber.cuh | 4 +- cpp/src/cugraph.cu | 71 +------------------- cpp/src/link_analysis/pagerank.cu | 2 +- cpp/src/link_prediction/jaccard.cu | 2 +- cpp/src/link_prediction/overlap.cu | 2 +- cpp/src/snmg/blas/spmv.cuh | 2 +- cpp/src/snmg/link_analysis/pagerank.cuh | 2 +- cpp/src/traversal/bfs.cu | 2 +- cpp/src/utilities/degree.cu | 86 +++++++++++++++++++++++++ cpp/src/{ => utilities}/graph_utils.cuh | 0 cpp/src/{ => utilities}/grmat.cu | 0 cpp/src/{ => utilities}/heap.cuh | 0 13 files changed, 97 insertions(+), 79 deletions(-) create mode 100644 cpp/src/utilities/degree.cu rename cpp/src/{ => utilities}/graph_utils.cuh (100%) rename cpp/src/{ => utilities}/grmat.cu (100%) rename cpp/src/{ => utilities}/heap.cuh (100%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index f357af04f9c..f0c68d181eb 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -229,7 +229,8 @@ link_directories("${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}" # CMAKE_CUDA_IMPLICIT ################################################################################################### # - library targets ------------------------------------------------------------------------------- add_library(cugraph SHARED - src/grmat.cu + src/utilities/grmat.cu + src/utilities/degree.cu src/cugraph.cu src/link_analysis/pagerank.cu src/traversal/bfs.cu diff --git a/cpp/src/converters/renumber.cuh b/cpp/src/converters/renumber.cuh index d1d0e9f03ba..5e2fa069267 100644 --- a/cpp/src/converters/renumber.cuh +++ b/cpp/src/converters/renumber.cuh @@ -34,8 +34,8 @@ #include #include "utilities/error_utils.h" -#include "graph_utils.cuh" -#include "heap.cuh" +#include "utilities/graph_utils.cuh" +#include "utilities/heap.cuh" #include "rmm_utils.h" namespace cugraph { diff --git a/cpp/src/cugraph.cu b/cpp/src/cugraph.cu index 505f84ed8d4..a5b1dd0e4ab 100644 --- a/cpp/src/cugraph.cu +++ b/cpp/src/cugraph.cu @@ -15,7 +15,7 @@ // Author: Alex Fender afender@nvidia.com #include -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "converters/COOtoCSR.cuh" #include "utilities/error_utils.h" #include "converters/renumber.cuh" @@ -250,75 +250,6 @@ gdf_error gdf_add_transposed_adj_list_impl (gdf_graph *graph) { return GDF_SUCCESS; } -gdf_error gdf_degree_impl(int n, int e, gdf_column* col_ptr, gdf_column* degree, bool offsets) { - if(offsets == true) { - dim3 nthreads, nblocks; - nthreads.x = min(n, CUDA_MAX_KERNEL_THREADS); - nthreads.y = 1; - nthreads.z = 1; - nblocks.x = min((n + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); - nblocks.y = 1; - nblocks.z = 1; - - switch (col_ptr->dtype) { - case GDF_INT32: cugraph::degree_offsets <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; - default: return GDF_UNSUPPORTED_DTYPE; - } - } - else { - dim3 nthreads, nblocks; - nthreads.x = min(e, CUDA_MAX_KERNEL_THREADS); - nthreads.y = 1; - nthreads.z = 1; - nblocks.x = min((e + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); - nblocks.y = 1; - nblocks.z = 1; - - switch (col_ptr->dtype) { - case GDF_INT32: cugraph::degree_coo <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; - default: return GDF_UNSUPPORTED_DTYPE; - } - } - return GDF_SUCCESS; -} - - -gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x) { - // Calculates the degree of all vertices of the graph - // x = 0: in+out degree - // x = 1: in-degree - // x = 2: out-degree - GDF_REQUIRE(graph->adjList != nullptr || graph->transposedAdjList != nullptr, GDF_INVALID_API_CALL); - int n; - int e; - if(graph->adjList != nullptr) { - n = graph->adjList->offsets->size -1; - e = graph->adjList->indices->size; - } - else { - n = graph->transposedAdjList->offsets->size - 1; - e = graph->transposedAdjList->indices->size; - } - - if(x!=1) { - // Computes out-degree for x=0 and x=2 - if(graph->adjList) - gdf_degree_impl(n, e, graph->adjList->offsets, degree, true); - else - gdf_degree_impl(n, e, graph->transposedAdjList->indices, degree, false); - } - - if(x!=2) { - // Computes in-degree for x=0 and x=1 - if(graph->adjList) - gdf_degree_impl(n, e, graph->adjList->indices, degree, false); - else - gdf_degree_impl(n, e, graph->transposedAdjList->offsets, degree, true); - } - return GDF_SUCCESS; -} - - gdf_error gdf_add_adj_list(gdf_graph *graph) { if (graph->adjList != nullptr) return GDF_SUCCESS; diff --git a/cpp/src/link_analysis/pagerank.cu b/cpp/src/link_analysis/pagerank.cu index 06f24c9b1f2..1943ba9f22b 100644 --- a/cpp/src/link_analysis/pagerank.cu +++ b/cpp/src/link_analysis/pagerank.cu @@ -25,7 +25,7 @@ #include -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "utilities/error_utils.h" #include diff --git a/cpp/src/link_prediction/jaccard.cu b/cpp/src/link_prediction/jaccard.cu index 56d62b4f6b3..fe3502e4356 100644 --- a/cpp/src/link_prediction/jaccard.cu +++ b/cpp/src/link_prediction/jaccard.cu @@ -19,7 +19,7 @@ * @file jaccard.cu * ---------------------------------------------------------------------------**/ -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "cugraph.h" #include "rmm_utils.h" #include "utilities/error_utils.h" diff --git a/cpp/src/link_prediction/overlap.cu b/cpp/src/link_prediction/overlap.cu index e6985ecda9e..cce0ac99752 100644 --- a/cpp/src/link_prediction/overlap.cu +++ b/cpp/src/link_prediction/overlap.cu @@ -19,7 +19,7 @@ * @file jaccard.cu * ---------------------------------------------------------------------------**/ -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "cugraph.h" #include "rmm_utils.h" #include "utilities/error_utils.h" diff --git a/cpp/src/snmg/blas/spmv.cuh b/cpp/src/snmg/blas/spmv.cuh index 27e15a1b1e3..8b7120a8e65 100644 --- a/cpp/src/snmg/blas/spmv.cuh +++ b/cpp/src/snmg/blas/spmv.cuh @@ -20,7 +20,7 @@ #pragma once #include "cub/cub.cuh" #include -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "snmg/utils.cuh" //#define SNMG_DEBUG diff --git a/cpp/src/snmg/link_analysis/pagerank.cuh b/cpp/src/snmg/link_analysis/pagerank.cuh index ea8637a81ac..6bf527bb4c0 100644 --- a/cpp/src/snmg/link_analysis/pagerank.cuh +++ b/cpp/src/snmg/link_analysis/pagerank.cuh @@ -20,7 +20,7 @@ #pragma once #include "cub/cub.cuh" #include -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "snmg/utils.cuh" #include "snmg/blas/spmv.cuh" //#define SNMG_DEBUG diff --git a/cpp/src/traversal/bfs.cu b/cpp/src/traversal/bfs.cu index 3018b08fcf3..c42be78943c 100644 --- a/cpp/src/traversal/bfs.cu +++ b/cpp/src/traversal/bfs.cu @@ -16,7 +16,7 @@ #include #include "rmm_utils.h" -#include "graph_utils.cuh" +#include "utilities/graph_utils.cuh" #include "bfs_kernels.cuh" using namespace bfs_kernels; diff --git a/cpp/src/utilities/degree.cu b/cpp/src/utilities/degree.cu new file mode 100644 index 00000000000..103330672f5 --- /dev/null +++ b/cpp/src/utilities/degree.cu @@ -0,0 +1,86 @@ +/* + * Copyright (c) 2019, 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 +#include "utilities/error_utils.h" +#include "utilities/graph_utils.cuh" + +gdf_error gdf_degree_impl(int n, int e, gdf_column* col_ptr, gdf_column* degree, bool offsets) { + if(offsets == true) { + dim3 nthreads, nblocks; + nthreads.x = min(n, CUDA_MAX_KERNEL_THREADS); + nthreads.y = 1; + nthreads.z = 1; + nblocks.x = min((n + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); + nblocks.y = 1; + nblocks.z = 1; + + switch (col_ptr->dtype) { + case GDF_INT32: cugraph::degree_offsets <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; + default: return GDF_UNSUPPORTED_DTYPE; + } + } + else { + dim3 nthreads, nblocks; + nthreads.x = min(e, CUDA_MAX_KERNEL_THREADS); + nthreads.y = 1; + nthreads.z = 1; + nblocks.x = min((e + nthreads.x - 1) / nthreads.x, CUDA_MAX_BLOCKS); + nblocks.y = 1; + nblocks.z = 1; + + switch (col_ptr->dtype) { + case GDF_INT32: cugraph::degree_coo <<>>(n, e, static_cast(col_ptr->data), static_cast(degree->data));break; + default: return GDF_UNSUPPORTED_DTYPE; + } + } + return GDF_SUCCESS; +} + + +gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x) { + // Calculates the degree of all vertices of the graph + // x = 0: in+out degree + // x = 1: in-degree + // x = 2: out-degree + GDF_REQUIRE(graph->adjList != nullptr || graph->transposedAdjList != nullptr, GDF_INVALID_API_CALL); + int n; + int e; + if(graph->adjList != nullptr) { + n = graph->adjList->offsets->size -1; + e = graph->adjList->indices->size; + } + else { + n = graph->transposedAdjList->offsets->size - 1; + e = graph->transposedAdjList->indices->size; + } + + if(x!=1) { + // Computes out-degree for x=0 and x=2 + if(graph->adjList) + gdf_degree_impl(n, e, graph->adjList->offsets, degree, true); + else + gdf_degree_impl(n, e, graph->transposedAdjList->indices, degree, false); + } + + if(x!=2) { + // Computes in-degree for x=0 and x=1 + if(graph->adjList) + gdf_degree_impl(n, e, graph->adjList->indices, degree, false); + else + gdf_degree_impl(n, e, graph->transposedAdjList->offsets, degree, true); + } + return GDF_SUCCESS; +} diff --git a/cpp/src/graph_utils.cuh b/cpp/src/utilities/graph_utils.cuh similarity index 100% rename from cpp/src/graph_utils.cuh rename to cpp/src/utilities/graph_utils.cuh diff --git a/cpp/src/grmat.cu b/cpp/src/utilities/grmat.cu similarity index 100% rename from cpp/src/grmat.cu rename to cpp/src/utilities/grmat.cu diff --git a/cpp/src/heap.cuh b/cpp/src/utilities/heap.cuh similarity index 100% rename from cpp/src/heap.cuh rename to cpp/src/utilities/heap.cuh From c5b5da935060cc872495d58bd24baed7f1cac99e Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 21:20:26 -0700 Subject: [PATCH 15/16] Reorg structure --- cpp/CMakeLists.txt | 2 +- cpp/src/{ => structure}/cugraph.cu | 0 2 files changed, 1 insertion(+), 1 deletion(-) rename cpp/src/{ => structure}/cugraph.cu (100%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index f0c68d181eb..bbd596551dc 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -231,7 +231,7 @@ link_directories("${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}" # CMAKE_CUDA_IMPLICIT add_library(cugraph SHARED src/utilities/grmat.cu src/utilities/degree.cu - src/cugraph.cu + src/structure/cugraph.cu src/link_analysis/pagerank.cu src/traversal/bfs.cu src/link_prediction/jaccard.cu diff --git a/cpp/src/cugraph.cu b/cpp/src/structure/cugraph.cu similarity index 100% rename from cpp/src/cugraph.cu rename to cpp/src/structure/cugraph.cu From 82ce0e5bf3cd2536e0a55172a57dd9078d2cfb93 Mon Sep 17 00:00:00 2001 From: Kumar Aatish Date: Mon, 13 May 2019 21:25:44 -0700 Subject: [PATCH 16/16] Fixed CHANGELOG --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index c200c2c0173..b5c0ac3b592 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -5,6 +5,7 @@ ## Improvements - PR #291 nvGraph is updated to use RMM instead of directly invoking cnmem functions. +- PR #286 Reorganized cugraph source directory ## Bug Fixes