Skip to content

Commit

Permalink
Merge pull request #5 from rapidsai/branch-0.8
Browse files Browse the repository at this point in the history
Latest updates
  • Loading branch information
aschaffer authored Jun 11, 2019
2 parents eb59faf + d70e3ab commit 788d727
Show file tree
Hide file tree
Showing 6 changed files with 59 additions and 62 deletions.
5 changes: 3 additions & 2 deletions CHANGELOG.md
Original file line number Diff line number Diff line change
@@ -1,8 +1,8 @@
# cuGraph 0.8.0 (Date TBD)

## New Features
- PR #287 SNMG power iteration step1
- PR #297 SNMG degree calculation
- PR #287 SNMG power iteration step1
- PR #297 SNMG degree calculation
- PR #300 Personalized Page Rank
- PR #302 SNMG CSR Pagerank (cuda/C++)
- PR #315 Weakly Connected Components adapted from cuML (cuda/C++)
Expand All @@ -22,6 +22,7 @@
- PR #311 Fixed bug in SNMG degree causing failure for three gpus
- PR #309 Update conda build recipes
- PR #314 Added datasets to gitignore
- PR #322 Updates to accommodate new cudf include file locations

# cuGraph 0.7.0 (10 May 2019)

Expand Down
4 changes: 1 addition & 3 deletions cpp/include/cugraph.h
Original file line number Diff line number Diff line change
Expand Up @@ -19,12 +19,10 @@
#include <cstdlib>
#include <cstdint>

#include <cudf.h>
#include <cudf/cudf.h>

#include "types.h"

#include "functions.h"

#include "algorithms.h"


16 changes: 8 additions & 8 deletions cpp/src/converters/renumber.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@

#include <thrust/scan.h>
#include <thrust/binary_search.h>
#include <cudf.h>
#include <cudf/cudf.h>
#include <cuda_runtime_api.h>

#include "utilities/error_utils.h"
Expand All @@ -44,7 +44,7 @@ namespace cugraph {
typedef uint32_t hash_type;
//typedef unsigned long long index_type;
typedef uint32_t index_type;

template <typename VertexIdType>
class HashFunctionObject {
public:
Expand Down Expand Up @@ -236,13 +236,13 @@ namespace cugraph {
[hash_bins_start, hash] __device__ (T_in vid) {
atomicAdd(hash_bins_start + hash(vid), detail::index_type{1});
});

thrust::for_each(rmm::exec_policy(stream)->on(stream),
dst, dst + size,
[hash_bins_start, hash] __device__ (T_in vid) {
atomicAdd(hash_bins_start + hash(vid), detail::index_type{1});
});


//
// Compute exclusive sum and copy it into both hash_bins_start and
Expand All @@ -261,15 +261,15 @@ namespace cugraph {
detail::index_type hash_offset = atomicAdd(&hash_bins_end[hash_index], 1);
hash_data[hash_offset] = vid;
});

thrust::for_each(rmm::exec_policy(stream)->on(stream),
dst, dst + size,
[hash_bins_end, hash_data, hash] __device__ (T_in vid) {
uint32_t hash_index = hash(vid);
detail::index_type hash_offset = atomicAdd(&hash_bins_end[hash_index], 1);
hash_data[hash_offset] = vid;
});

//
// Now we need to dedupe the hash bins
//
Expand Down Expand Up @@ -311,7 +311,7 @@ namespace cugraph {
*new_size = temp;

ALLOC_TRY(numbering_map, (*new_size) * sizeof(T_in), nullptr);

T_in * local_numbering_map = *numbering_map;

thrust::for_each(rmm::exec_policy(stream)->on(stream),
Expand All @@ -332,7 +332,7 @@ namespace cugraph {

return GDF_SUCCESS;
}

}

#endif
24 changes: 11 additions & 13 deletions cpp/src/tests/components/con_comp_test.cu
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@

#include "gtest/gtest.h"
#include "high_res_clock.h"
#include <cudf.h>
#include <cudf/cudf.h>
#include "cuda_profiler_api.h"

#include <cugraph.h>
Expand Down Expand Up @@ -50,7 +50,7 @@ namespace{ //un-nammed
{
matrix_file = rhs.matrix_file;
}

Usecase& operator = (const Usecase& rhs)
{
matrix_file = rhs.matrix_file;
Expand All @@ -64,19 +64,19 @@ namespace{ //un-nammed
private:
std::string matrix_file;
};

}//end un-nammed namespace

struct Tests_Weakly_CC : ::testing::TestWithParam<Usecase>
{
Tests_Weakly_CC() { }
static void SetupTestCase() { }
static void TearDownTestCase() {
static void TearDownTestCase() {
if (PERF) {
for (unsigned int i = 0; i < weakly_cc_time.size(); ++i) {
std::cout << weakly_cc_time[i] << std::endl;
}
}
}
}
virtual void SetUp() { }
virtual void TearDown() { }
Expand All @@ -85,13 +85,13 @@ struct Tests_Weakly_CC : ::testing::TestWithParam<Usecase>

void run_current_test(const Usecase& param) {
const ::testing::TestInfo* const test_info =::testing::UnitTest::GetInstance()->current_test_info();
std::stringstream ss;
std::stringstream ss;
std::string test_id = std::string(test_info->test_case_name()) + std::string(".") + std::string(test_info->name()) + std::string("_") + getFileName(param.get_matrix_file())+ std::string("_") + ss.str().c_str();
cudaStream_t stream{nullptr};

int m, k, nnz; //
MM_typecode mc;

HighResClock hr_clock;
double time_tmp;

Expand All @@ -102,7 +102,7 @@ struct Tests_Weakly_CC : ::testing::TestWithParam<Usecase>
ASSERT_TRUE(mm_is_matrix(mc));
ASSERT_TRUE(mm_is_coordinate(mc));
ASSERT_TRUE(mm_is_symmetric(mc));//weakly cc only works w/ undirected graphs, for now;

// Allocate memory on host
std::vector<int> cooRowInd(nnz);
std::vector<int> cooColInd(nnz);
Expand Down Expand Up @@ -144,18 +144,18 @@ struct Tests_Weakly_CC : ::testing::TestWithParam<Usecase>
cudaDeviceSynchronize();
}
EXPECT_EQ(status,GDF_SUCCESS);

}
};

std::vector<double> Tests_Weakly_CC::weakly_cc_time;

TEST_P(Tests_Weakly_CC, Weakly_CC) {
run_current_test(GetParam());
}

// --gtest_filter=*simple_test*
INSTANTIATE_TEST_CASE_P(simple_test, Tests_Weakly_CC,
INSTANTIATE_TEST_CASE_P(simple_test, Tests_Weakly_CC,
::testing::Values( Usecase("networks/dolphins.mtx")
//Usecase("networks/coPapersDBLP.mtx"),
//Usecase("networks/coPapersCiteseer.mtx"),
Expand All @@ -173,5 +173,3 @@ int main(int argc, char **argv) {

return RUN_ALL_TESTS();
}


66 changes: 33 additions & 33 deletions cpp/src/utilities/grmat.cu
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@

#include <gunrock/util/shared_utils.cuh>

#include <cudf.h>
#include <cudf/cudf.h>
#include <thrust/extrema.h>
#include "utilities/error_utils.h"
#include "graph_utils.cuh"
Expand All @@ -44,7 +44,7 @@ using namespace gunrock::util;
using namespace gunrock::graphio;
using namespace gunrock::graphio::grmat;

template <typename VertexId, typename Value, typename SizeT>
template <typename VertexId, typename Value, typename SizeT>
__global__ void Remove_Self_Loops (VertexId* row, VertexId* col, Value* val, SizeT edges)
{
SizeT i = (SizeT)blockIdx.x * blockDim.x + threadIdx.x;
Expand All @@ -54,7 +54,7 @@ __global__ void Remove_Self_Loops (VertexId* row, VertexId* col, Value* val, Siz
if (row[i] == col[i])
{
col[i] = 0;
}
}
}
}

Expand Down Expand Up @@ -98,9 +98,9 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine
bool quiet = false;

typedef Coo_nv<VertexId, Value> EdgeTupleType;

cpu_timer.Start();

if (args->CheckCmdLineFlag ("rmat_scale") && args->CheckCmdLineFlag ("rmat_nodes"))
{
printf ("Please mention scale or nodes, not both \n");
Expand All @@ -111,8 +111,8 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine
printf ("Please mention edgefactor or edge, not both \n");
return GDF_UNSUPPORTED_METHOD;
}
self_loops = args->CheckCmdLineFlag ("rmat_self_loops");

self_loops = args->CheckCmdLineFlag ("rmat_self_loops");
// graph construction or generation related parameters
if (args -> CheckCmdLineFlag("normalized"))
undirected = args -> CheckCmdLineFlag("rmat_undirected");
Expand All @@ -138,10 +138,10 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine
rmat_seed = -1;
}
EdgeTupleType coo;

if (undirected == true)
{
rmat_all_edges = 2 * rmat_edges;
rmat_all_edges = 2 * rmat_edges;
}
else
{
Expand Down Expand Up @@ -169,8 +169,8 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine
printf ("---------Graph properties-------\n"
" Undirected : %s\n"
" Nodes : %lld\n"
" Edges : %lld\n"
" a = %f, b = %f, c = %f, d = %f\n\n\n", ((undirected == true)? "True": "False"), (long long)rmat_nodes,
" Edges : %lld\n"
" a = %f, b = %f, c = %f, d = %f\n\n\n", ((undirected == true)? "True": "False"), (long long)rmat_nodes,
(long long)(rmat_edges * ((undirected == true)? 2: 1)), rmat_a, rmat_b, rmat_c, rmat_d);
}

Expand All @@ -197,8 +197,8 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine
cpu_timer2.Start();
cudaError_t status = cudaSuccess;
if(val == nullptr)
status = BuildRmatGraph_coo_nv<false, VertexId, SizeT, Value, EdgeTupleType>(rmat_nodes, rmat_edges, coo, undirected,
rmat_a, rmat_b, rmat_c, rmat_d, rmat_vmultipiler, rmat_vmin, rmat_seed,
status = BuildRmatGraph_coo_nv<false, VertexId, SizeT, Value, EdgeTupleType>(rmat_nodes, rmat_edges, coo, undirected,
rmat_a, rmat_b, rmat_c, rmat_d, rmat_vmultipiler, rmat_vmin, rmat_seed,
quiet, temp_devices.size(), gpu_idx);
else
status = BuildRmatGraph_coo_nv<true, VertexId, SizeT, Value, EdgeTupleType>(rmat_nodes, rmat_edges, coo, undirected,
Expand All @@ -222,15 +222,15 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine

return GDF_CUDA_ERROR;
}

int block_size = (sizeof(VertexId) == 4) ? 1024 : 512;
int grid_size = rmat_all_edges / block_size + 1;

if (util::SetDevice(gpu_idx[0]))
return GDF_CUDA_ERROR;
if ((self_loops != false) && (val != nullptr))
{
Remove_Self_Loops
{
Remove_Self_Loops
<VertexId, Value, SizeT>
<<<grid_size, block_size, 0>>>
(coo.row, coo.col, coo.val, rmat_all_edges);
Expand All @@ -242,26 +242,26 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine

VertexId nodes_row = 0;
VertexId nodes_col = 0;

cudaMemcpy((void*)&nodes_row, (void*)&(coo.row[rmat_all_edges-1]), sizeof(VertexId), cudaMemcpyDeviceToHost);

tmp = thrust::max_element(rmm::exec_policy(stream)->on(stream),
thrust::device_pointer_cast((VertexId*)(coo.col)),
thrust::device_pointer_cast((VertexId*)(coo.col)),
thrust::device_pointer_cast((VertexId*)(coo.col + rmat_all_edges)));
nodes_col = tmp[0];

VertexId max_nodes = (nodes_row > nodes_col)? nodes_row: nodes_col;

cpu_timer.Stop();

if ((src != nullptr) && (dest != nullptr))
{
src->data = coo.row;
{
src->data = coo.row;
src->size = rmat_all_edges;
src->valid = nullptr;
dest->data = coo.col;
dest->size = rmat_all_edges;
src->valid = nullptr;

dest->data = coo.col;
dest->size = rmat_all_edges;
dest->valid = nullptr;
}
else
Expand All @@ -274,20 +274,20 @@ gdf_error main_(gdf_column *src, gdf_column *dest, gdf_column *val, CommandLine
ALLOC_FREE_TRY(coo.val, stream);
if (!quiet)
printf ("Error : Pointers for gdf column are null, releasing allocated memory for graph\n");

return GDF_CUDA_ERROR;
}

if (val != nullptr)
{
val->data = coo.val;
val->size = rmat_all_edges;
val->data = coo.val;
val->size = rmat_all_edges;
val->valid = nullptr;
}

vertices = max_nodes+1;
edges = rmat_all_edges;

if (!quiet)
printf ("Time to generate the graph %f ms\n"
"Total time %f ms\n", cpu_timer2.ElapsedMillis(), cpu_timer.ElapsedMillis());
Expand Down Expand Up @@ -327,7 +327,7 @@ gdf_error gdf_grmat_gen (const char* argv, size_t& vertices, size_t& edges, gdf_
if (src == nullptr || dest == nullptr)
{
free_args(argc, arg);
return GDF_DATASET_EMPTY;
return GDF_DATASET_EMPTY;
}

GDF_REQUIRE ((src->dtype == dest->dtype), GDF_DTYPE_MISMATCH);
Expand Down
6 changes: 3 additions & 3 deletions python/cugraph/graph/c_graph.pxd
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
from libcpp cimport bool

cdef extern from "cudf.h":
cdef extern from "cudf/cudf.h":

ctypedef size_t gdf_size_type

Expand All @@ -16,7 +16,7 @@ cdef extern from "cudf.h":
ctypedef struct gdf_dtype_extra_info:
gdf_time_unit time_unit

ctypedef enum gdf_error:
ctypedef enum gdf_error:

pass

Expand Down Expand Up @@ -99,4 +99,4 @@ cdef extern from "cugraph.h":
cdef gdf_error gdf_add_transposed_adj_list(gdf_graph *graph)
cdef gdf_error gdf_delete_transposed_adj_list(gdf_graph *graph)

cdef gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x)
cdef gdf_error gdf_degree(gdf_graph *graph, gdf_column *degree, int x)

0 comments on commit 788d727

Please sign in to comment.