gunrock / gunrock Goto Github PK

View Code? Open in Web Editor NEW

940.0 77.0 198.0 76.38 MB

Programmable CUDA/C++ GPU Graph Analytics

Home Page: https://gunrock.github.io/gunrock/

License: Apache License 2.0

CMake 2.92% Cuda 16.59% Makefile 1.23% C++ 78.62% C 0.02% Shell 0.62%

gunrock cuda graph-processing graph-analytics gpu graph-primitives graph-engine hpc cxx algorithms

gunrock's Introduction

Gunrock: CUDA/C++ GPU Graph Analytics

Examples	Project Template	Documentation	GitHub Actions

Gunrock¹ is a CUDA library for graph-processing designed specifically for the GPU. It uses a high-level, bulk-synchronous/asynchronous, data-centric abstraction focused on operations on vertex or edge frontiers. Gunrock achieves a balance between performance and expressiveness by coupling high-performance GPU computing primitives and optimization strategies, particularly in the area of fine-grained load balancing, with a high-level programming model that allows programmers to quickly develop new graph primitives that scale from one to many GPUs on a node with small code size and minimal GPU programming knowledge.

Branch	Purpose	Version	Status
`main`	Default branch, ported from `gunrock/essentials`, serves as the official release branch.	$\geq$ `2.x.x`	Active
`develop`	Development feature branch, ported from `gunrock/essentials`.	$\geq$ `2.x.x`	Active
`master`	Previous release branch for `gunrock/gunrock` version `1.x.x` interface, preserves all commit history.	$\leq$ `1.x.x`	Deprecated
`dev`	Previous development branch for `gunrock/gunrock`. All changes now merged in `master`.	$\leq$ `1.x.x`	Deprecated

Quick Start Guide

Before building Gunrock make sure you have CUDA Toolkit² installed on your system. Other external dependencies such as NVIDIA/thrust, NVIDIA/cub, etc. are automatically fetched using cmake.

git clone https://github.com/gunrock/gunrock.git
cd gunrock
mkdir build && cd build
cmake .. 
make sssp # or for all algorithms, use: make -j$(nproc)
bin/sssp ../datasets/chesapeake/chesapeake.mtx

Implementing Graph Algorithms

For a detailed explanation, please see the full documentation. The following example shows simple APIs using Gunrock's data-centric, bulk-synchronous programming model, we implement Breadth-First Search on GPUs. This example skips the setup phase of creating a problem_t and enactor_t struct and jumps straight into the actual algorithm.

We first prepare our frontier with the initial source vertex to begin push-based BFS traversal. A simple f->push_back(source) places the initial vertex we will use for our first iteration.

void prepare_frontier(frontier_t* f,
                      gcuda::multi_context_t& context) override {
  auto P = this->get_problem();
  f->push_back(P->param.single_source);
}

We then begin our iterative loop, which iterates until a convergence condition has been met. If no condition has been specified, the loop converges when the frontier is empty.

void loop(gcuda::multi_context_t& context) override {
  auto E = this->get_enactor();   // Pointer to enactor interface.
  auto P = this->get_problem();   // Pointer to problem (data) interface.
  auto G = P->get_graph();        // Graph that we are processing.

  auto single_source = P->param.single_source;  // Initial source node.
  auto distances = P->result.distances;         // Distances array for BFS.
  auto visited = P->visited.data().get();       // Visited map.
  auto iteration = this->iteration;             // Iteration we are on.

  // Following lambda expression is applied on every source,
  // neighbor, edge, weight tuple during the traversal.
  // Our intent here is to find and update the minimum distance when found.
  // And return which neighbor goes in the output frontier after traversal.
  auto search = [=] __host__ __device__(
                      vertex_t const& source,    // ... source
                      vertex_t const& neighbor,  // neighbor
                      edge_t const& edge,        // edge
                      weight_t const& weight     // weight (tuple).
                      ) -> bool {
    auto old_distance =
      math::atomic::min(&distances[neighbor], iteration + 1);
    return (iteration + 1 < old_distance);
  };

  // Execute advance operator on the search lambda expression.
  // Uses load_balance_t::block_mapped algorithm (try others for perf. tuning.)
  operators::advance::execute<operators::load_balance_t::block_mapped>(
    G, E, search, context);
}

include/gunrock/algorithms/bfs.hxx

How to Cite Gunrock & Essentials

Thank you for citing our work.

@article{Wang:2017:GGG,
  author =	 {Yangzihao Wang and Yuechao Pan and Andrew Davidson
                  and Yuduo Wu and Carl Yang and Leyuan Wang and
                  Muhammad Osama and Chenshan Yuan and Weitang Liu and
                  Andy T. Riffel and John D. Owens},
  title =	 {{G}unrock: {GPU} Graph Analytics},
  journal =	 {ACM Transactions on Parallel Computing},
  year =	 2017,
  volume =	 4,
  number =	 1,
  month =	 aug,
  pages =	 {3:1--3:49},
  doi =		 {10.1145/3108140},
  ee =		 {http://arxiv.org/abs/1701.01170},
  acmauthorize = {https://dl.acm.org/doi/10.1145/3108140?cid=81100458295},
  url =		 {http://escholarship.org/uc/item/9gj6r1dj},
  code =	 {https://github.com/gunrock/gunrock},
  ucdcite =	 {a115},
}

@InProceedings{Osama:2022:EOP,
  author =	 {Muhammad Osama and Serban D. Porumbescu and John D. Owens},
  title =	 {Essentials of Parallel Graph Analytics},
  booktitle =	 {Proceedings of the Workshop on Graphs,
                  Architectures, Programming, and Learning},
  year =	 2022,
  series =	 {GrAPL 2022},
  month =	 may,
  pages =	 {314--317},
  doi =		 {10.1109/IPDPSW55747.2022.00061},
  url =          {https://escholarship.org/uc/item/2p19z28q},
}

Copyright & License

Gunrock is copyright The Regents of the University of California. The library, examples, and all source code are released under Apache 2.0.

This repository has been moved from https://github.com/gunrock/essentials and the previous history is preserved with tags and under master branch. Read more about gunrock and essentials in our vision paper: Essentials of Parallel Graph Analytics. ↩
Recommended CUDA v11.5.1 or higher due to support for stream ordered memory allocators. ↩

gunrock's People

Contributors

Stargazers

Watchers

Forkers

wangshaohua bldrake yzhwang aglenis 1duo shenguojun wsumeilin anshumang sungsoo jayavanth vvishal damien- kongfujianlong laurawly welcheb kalatestimine tkelman pavanky prayagverma zcourts devavret neoblizz ffarhour shariiiyy ptheywood lykkos cequencer codeaudit tigran-hakobyan nsakharnykh graphprocessor gumi-presentation-by-dzh rtvt123 slimbenamor alexanderjpowell asherliu davidtranno1 archenroot arokem guohaoqiang gmolini ml-lab foxspy nagyistge desert0616 adamharries iraikov tonythomascn limin2021 rchandramouli tass-lab teju85 sepandhaghighi alexperrone j143-zz xyuan seungbinsong vishalbelsare jkelling caomw vcave solertis aicampbell shivampdesai sgpyc gabrielcc2 bharathbs93 strikew bookug ctcyang wang-lp shigangli gwcloudlab marziehlenjani zhuoyw hooman96 faustineinsun ank845 aliakbr junliangma siegfriedgreg xiezhq-hermann rozgo zlcook bkj pwrliang owensgroup agalup ajaybrahmakshatriya iroy30 chomolungma rgsl888prabhu wangxiaoyunanne porumbes wennyyustalim brownproton allenhsin amusaafir pstutz yyht

gunrock's Issues

Support for Nvidia Tesla C2075 Fermi architecture

Hi,

I have installed gunrock-v0.2-145-g57842ba on Nvidia Tesla C2075, while running the test: $GUNROCK_ROOT/tests/bfs after make, it gets hanged although process consumes memory on gpu seen from "nvidia-smi". May I know whether this particular GPU which has Fermi arch with compute capability 2.0 supports this version of gunrock or not? If not, is there any way to make gunrock running for this architecture? Or is there any other version of gunrock which is compatible to this architecture?

gitsha1.c: should not be in repo

I'm not 100% sure I did this right, but I think I did: gitsha1.c is generated by CMake and should not be in the repo. However, in the merge, it is in the repo and I believe it should not be. (The header and .in files SHOULD be in the repo, and are.) Whoever added it to the repo, please check that you did the right thing, and remove it if it shouldn't be there.

http://stackoverflow.com/questions/1435953/how-can-i-pass-git-sha1-to-compiler-as-definition-using-cmake

Improve Gunrock's Documentation

Need to add details on how to create new graph primitives with problem, functor, and enactor. Proposed method: Add functionality and interface explanation for app/template (the empty sample template graph primitive) so that users can create a minimal running graph primitive by walking through the template code.

To change openmp routines to use C++11 threads

getting clang working with openmp is too much to just let gunrock run on Mac OS X. Need to use C++11 multi-thread support for the graph generators. But C++11 mutex support is still not available (or at least not though the usual way) under current version of clang coming with Mac OS X (Apple LLVM version 6.1.0, clang-602.0.53, based on LLVM 3.6.0svn), need to wait till it's supported, and mgpu-merge branch should be able to compile under Mac OS X.

refine documentation creation, document it

I'm a little unclear on how to generate documentation. I understand we can use Jekyll, which I haven't used. But our docs are split between three places: (1) the gunrock-docs repo; (2) the README in the gunrock repo (which produces nice HTML on gunrock.github.io); and (3) the "doc" directory, which also generates the doxygen documentation (but does not appear to produce nice HTML on gunrock.github.io).

The proximate reason for this request is that I want to start putting Vega graphs into our docs, and that means I really need to understand how to put them in since they need JS and HTML hackery, and I don't know how/where the source files ought to go.

I'd like to see/understand:

A brief set of internal instructions on how/where to put docs. ("If you want to add a page to the documentation, here's where you put it, and here's how you make sure it's included in the {doxygen, gunrock.github.io's HTML} produced docs."
Non-hardcoded links for the web-hosted docs. Currently we have .md files which should render nicely as HTML files, but we link to them directly into the doc repo with a hardcoded version number (e.g., http://gunrock.github.io/gunrock/doc/0.3/programming_model.html). That should definitely be cleaner.
Jekyll. Should we be using this?

Assigning YZH simply because I think he built the initial bits here, but anyone can chime in.

merge multi-dev branch to master todos

TODO: Merge mgpu-dev branch into master. Temporary branch for merging: mgpu-merge
[Link: https://github.com/gunrock/gunrock/tree/mgpu-merge]

~~Solve merge conflict after commit 15cc802;~~
- merged in a95cd42
  BC, BFS, CC, HITS, PR, SSSP, TOPK, WTF passed tests
- ~~DOBFS yield INCORRECT results on directed inputs~~ fixed in 2d657ce
- SALSA has CUDA errors on large datasets
  (cudaMalloc error 77, [../../gunrock/util/array_utils.cuh, 182 @ gpu 0] values cudaMalloc failed (CUDA error 2: out of memory) on Kron datasets, maybe related to #25)
Align MIS, ~~MST, template, VIS,~~ to new interfaces;
- template done in 96bf635
- VIS done in 592fe0c
- MST finished in 460afc6
~~CMAKE and shared library build;~~
- cmake fixed BC, BFS, CC, HITS, PR, SSSP, DOBFS, SALSA, TOPK, WTF in fe44596
- cmake fixed for MST in 50955b7
- shared library interface for BFS fixed in 9a98f97
- shared library interface for BC fixed in 62051d2
- shared library interface for CC fixed in 7533461
- shared library interface for PageRank fixed in ad85d48
- shared library interface for SSSP fixed in 9c00efa
- cmake test fixed in 052e878
- py interface working in 51bdc51, all done here
Documentations;
Code comments;

Long running time for BGL betweenness centrality code on huge datasets

The CPU validation code for betweenness centrality using Boost Graph Library will take a long time on huge datasets. For example, for dataset ak2010, it will take about half an hour on a machine with Intel Core i5 CPU and 4G memory. Larger dataset such as socLiveJournal1 would probably take days to finish. We need to add precomputed validation results for BC as binary files to reduce the testing time.

simple_example compiler warning [line 526] on string literal to char * conversion

[ 71%] Building NVCC (Device) object simple_example/CMakeFiles/simple_example.dir//./simple_example_generated_simple_example.cu.o
/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu(526): warning: conversion from a string literal to "char *" is deprecated
          detected during instantiation of "void RunTests(gunrock::Csr<VertexId, Value, SizeT> &, gunrock::util::CommandLineArgs &, mgpu::CudaContext &) [with VertexId=int, Value=float, SizeT=int]"
(779): here

/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu(526): warning: conversion from a string literal to "char *" is deprecated

duplicate symbols in OS X build

Just wondering if this has come up in any other builds.

[ 18%] Building CXX object gunrock/CMakeFiles/gunrock.dir/__/externals/moderngpu/src/mgpuutil.cpp.o
Linking CXX shared library ../lib/libgunrock.dylib
duplicate symbol __ZN7gunrock4util12stringprintfEPKcz in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bfs/gunrock_generated_bfs_app.cu.o
duplicate symbol __ZN7gunrock7graphio5SprngEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bfs/gunrock_generated_bfs_app.cu.o
duplicate symbol __ZN7gunrock7graphio4FlipEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bfs/gunrock_generated_bfs_app.cu.o
duplicate symbol __ZN7gunrock7graphio10VaryParamsEPdS1_S1_S1_ in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bfs/gunrock_generated_bfs_app.cu.o
duplicate symbol __ZN7gunrock4util12stringprintfEPKcz in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bc/gunrock_generated_bc_app.cu.o
duplicate symbol __ZN7gunrock7graphio5SprngEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bc/gunrock_generated_bc_app.cu.o
duplicate symbol __ZN7gunrock7graphio4FlipEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bc/gunrock_generated_bc_app.cu.o
duplicate symbol __ZN7gunrock7graphio10VaryParamsEPdS1_S1_S1_ in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/bc/gunrock_generated_bc_app.cu.o
duplicate symbol __ZN7gunrock4util12stringprintfEPKcz in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/cc/gunrock_generated_cc_app.cu.o
duplicate symbol __ZN7gunrock7graphio5SprngEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/cc/gunrock_generated_cc_app.cu.o
duplicate symbol __ZN7gunrock7graphio4FlipEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/cc/gunrock_generated_cc_app.cu.o
duplicate symbol __ZN7gunrock7graphio10VaryParamsEPdS1_S1_S1_ in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/cc/gunrock_generated_cc_app.cu.o
duplicate symbol __ZN7gunrock4util12stringprintfEPKcz in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/pr/gunrock_generated_pr_app.cu.o
duplicate symbol __ZN7gunrock7graphio5SprngEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/pr/gunrock_generated_pr_app.cu.o
duplicate symbol __ZN7gunrock7graphio4FlipEv in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/pr/gunrock_generated_pr_app.cu.o
duplicate symbol __ZN7gunrock7graphio10VaryParamsEPdS1_S1_S1_ in:
    CMakeFiles/gunrock.dir/app/topk/gunrock_generated_topk_app.cu.o
    CMakeFiles/gunrock.dir/app/pr/gunrock_generated_pr_app.cu.o
ld: 16 duplicate symbols for architecture x86_64

priority_queue/kernel has a warning

/home/jowens/working/gunrock/gunrock/priority_queue/kernel.cuh(487): warning: variable "far_size" was set but never used

Compile error

error: no instance of function template "gunrock::oprtr::advance::LaunchKernel" matches the argument list
argument types are: (int *, gunrock::app::EnactorStats, gunrock::app::FrontierAttribute, gunrock::app::bfs::BFSProblem::DataSlice *, VertexId *, __nv_bool *, __nv_bool *, unsigned int *, int *, int *, VertexId *, int *, int *, int *, SizeT *, VertexId *, int, int, gunrock::util::CtaWorkProgressLifetime, mgpu::CudaContext, gunrock::oprtr::advance::TYPE)
detected during:
instantiation of "cudaError_t gunrock::app::bfs::BFSEnactor::EnactBFS(mgpu::CudaContext &, BFSProblem *, BFSProblem::VertexId, int) [with INSTRUMENT=true, AdvanceKernelPolicy=gunrock::oprtr::advance::KernelPolicy, 300, true, 8, 10, 8, 4096, 1, 0, 5, 32, 512, 7, LB>, FilterKernelPolicy=gunrock::oprtr::filter::KernelPolicy, 300, true, 0, true, 8, 8, 1, 0, 5, 5, 8>, BFSProblem=gunrock::app::bfs::BFSProblem]"
(495): here
instantiation of "cudaError_t gunrock::app::bfs::BFSEnactor::Enact(mgpu::CudaContext &, BFSProblem *, BFSProblem::VertexId, int, int) [with INSTRUMENT=true, BFSProblem=gunrock::app::bfs::BFSProblem]"

I am using:
CUDA 6.5,
GCC / G++ 4.4.7
CENTOS 6.4
Boost version: 1.57.0
cub 1.3.2
modernGPU

any / all conflicts / please only include the std:: things we need

@ydwu @yzhwang Can you narrow down the "using namespace std" directive? Import only the symbols we need, not the whole namespace?

https://isocpp.org/wiki/faq/coding-standards#using-namespace-std

e.g. "using std::cout" ?

Cliff says:

Hi John,

Further insight from our compiler team:

This ambiguity occurs because the user code introduces names from the std namespace in global scope with a "using" directive:
27 "/Users/jowens/Documents/working/gunrock/gunrock/csr.cuh" 2
using namespace std;

To avoid the ambiguity error, the user can explicitly qualify the use of __any and __all with scope resolution operator ("::").
i.e. old code:
599 "/Users/jowens/Documents/working/gunrock/externals/cub/cub/block/../util_ptx.cuh"
return __any(cond);

573 "/Users/jowens/Documents/working/gunrock/externals/cub/cub/block/../util_ptx.cuh"
return __all(cond);

code with fix:
599 "/Users/jowens/Documents/working/gunrock/externals/cub/cub/block/../util_ptx.cuh"
return ::__any(cond);

573 "/Users/jowens/Documents/working/gunrock/externals/cub/cub/block/../util_ptx.cuh"
return ::__all(cond);

Can you please try this out and let us know if that takes care of it? If so, I'll ask Duane to apply the same fix to CUB as well.

Thanks,
Cliff

Making pull-based traversal in dobfs a general operator.

New advance should look like this:
if (meets_some_metric)
push-based-advance //could be either TWC or LB
pull-based-advance //again, could be either TWC or LB

Need to be very careful when switching/generating frontiers for pull/push based advance.

Misalignment of BFS perf (traversal-mode=1) between single-node and master(multi-GPU) version

@sgpyc and @ydwu ,

I noticed that there is a misalignment for BFS perf when using traversal-mode=1 (dynamic grouping strategy of Merrill et al's paper) between the single-node branch and the master branch. It happens on two regular graphs: rgg_n24 and europe_osm. Could you guys look into this problem and find out why master branch performs worse when using the same traversal-mode=1 as the single-node branch? Thanks!

Better way to handle kernel call parameters

CUDA stores kernel parameters in shared memory. The limit is 256 bytes total. If using 64-bit, that's only 32 pointers/ints/floats. Currently Gunrock's number of arguments for kernel calls are getting more and more. Putting them all together in a struct and passing a handle would make the interface much cleaner.

cmake no longer has wizard mode (at least on OS X)

$ cmake -i ../gunrock
The "cmake -i" wizard mode is no longer supported.
Use the -D option to set cache values on the command line.
Use cmake-gui or ccmake for an interactive dialog.
$ cmake --version
cmake version 3.2.2

CMake suite maintained and supported by Kitware (kitware.com/cmake).

(docs need to be updated: http://gunrock.github.io/gunrock/doc/0.2/building_gunrock.html)

Who-To-Follow cudaMemcpy d_rank_curr[src] failed

[yuduo@mario]build$ bin/who_to_follow market /data/gunrock_dataset/small/chesapeake.mtx
Reading from /data/gunrock_dataset/small/chesapeake.mtx:
Parsing MARKET COO format (39 nodes, 170 directed edges)... Done parsing (0s).
Converting 39 vertices, 170 directed edges (unordered tuples) to CSR format...
Done converting (0s).

Degree Histogram (39 vertices, 170 edges):
Degree 0: 4 (10.26%)
Degree 2^0: 1 (2.56%)
Degree 2^1: 13 (33.33%)
Degree 2^2: 15 (38.46%)
Degree 2^3: 6 (15.38%)

PPR Time: 0.539072
CoT Time: 0.765632
[/home/yuduo/lab_project/gunrock/gunrock/app/wtf/wtf_enactor.cuh, 458] WTFProblem cudaMemcpy d_rank_curr[src] failed (CUDA error 11: invalid argument)
[/home/yuduo/lab_project/gunrock/tests/wtf/test_wtf.cu, 424] HITS Problem Enact Failed (CUDA error 11: invalid argument)

GPU result.Top 10 Page Ranks:
Vertex ID: 16, Page Rank: 0.000000
Vertex ID: 6, Page Rank: 0.000000
Vertex ID: 7, Page Rank: 0.000000
Vertex ID: 10, Page Rank: 0.000000
Vertex ID: 11, Page Rank: 0.000000
Vertex ID: 12, Page Rank: 0.000000
Vertex ID: 21, Page Rank: 0.000000
Vertex ID: 22, Page Rank: 0.000000
Vertex ID: 34, Page Rank: 0.000000
Vertex ID: 13, Page Rank: 0.000000
[GPU Who-To-Follow] finished. elapsed: 2.265 ms

Compile warning in topk_app.cu (VALUE_DOUBLE not handled in switch)

[ 50%] Building NVCC (Device) object gunrock/CMakeFiles/gunrock.dir/__/tests/topk/./gunrock_generated_topk_app.cu.o
/Users/jowens/Documents/working/gunrock/gunrock/../tests/topk/topk_app.cu:109:9: warning:
      enumeration value 'VALUE_DOUBLE' not handled in switch [-Wswitch]
switch (data_type.VALUE_TYPE)
        ^
1 warning generated.

Possible CUDA memory allocation failure on huge datasets

When running BFS (and possibly BC) on kron_g500-logn21 and soc-LiveJournal1 for undirected graph and choose to keep predecessor for each node, there is possible CUDA memory allocation failure if the device memory of the card you are using is not enough. Need to add free device memory detecting code.

BFS Algorithm Detail

Hi,
I am exploring the BFS algorithm implementation in cuda gunrock, in this regards I read the paper "High Performance and Scalable GPU Graph Traversal". However the paper lists a number of algorithms. Could you please let me know the exact algorithm that is used in gunrock. Where can I get more details about this algorithm?

Thanks

Build Error

Hi,
I'm trying to build using CMake and getting the following error
'Parse error. Expected "(", got newline with text "
".'

Gunrock implementation

I have gone through the Gunrock paper "Gunrock: A High-Performance Graph Processing Library on the GPU", but still unable to understand the implementation of Gunrock sssp algorithm. Can you please tell details about the implementation of cuda version of sssp(gunrock).

potential problem of running BFS on kron_g500_logn17 while enabling idempotence, src=0, undirected

./bin/test_bfs_6.5_x86_64 market /data/gunrock_dataset/large/kron_g500-logn17/kron_g500-logn17.mtx --src=0 --undirected --idempotence --iteration-num=10

[../../gunrock/util/cta_work_progress.cuh, 330] CtaWorkProgress cudaMemcpy d_counters failed (CUDA error 77: an illegal memory access was encountered)
[test_bfs.cu, 387] BFS Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bfs/bfs_problem.cuh, 156] BFSProblem cudaMemcpy d_labels failed (CUDA error 77: an illegal memory access was encountered)
[test_bfs.cu, 399] BFS Problem Data Extraction Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bfs/bfs_problem.cuh, 118] GpuSlice cudaFree d_labels failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bfs/bfs_problem.cuh, 120] GpuSlice cudaFree d_visited_mask failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bfs/bfs_problem.cuh, 121] GpuSlice cudaFree data_slices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 126] GpuSlice cudaFree d_row_offsets failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 127] GpuSlice cudaFree d_column_indices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bfs/bfs_enactor.cuh, 149] BFSEnactor cudaFreeHost done failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bfs/bfs_enactor.cuh, 152] BFSEnactor cudaEventDestroy throttle_event failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 120] EnactorBase cudaFree d_node_locks failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 121] EnactorBase cudaFree d_node_locks_out failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed: (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed: (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/cta_work_progress.cuh, 230] CtaWorkProgress cudaFree d_counters failed: (CUDA error 77: an illegal memory access was encountered)

clang issues with templates ("error: non-type template argument depends on a template parameter of the partial specialization")

These are clang vs. gcc errors. (Who knows what VC++ will do here.)
GCC likes them, clang does not. Mac uses clang by default (and I can
switch to gcc if I need to switch to gcc). This is not high priority
for a release, but we wanna figure out what's wrong, and if we can
write code that works on all compilers, that's certainly desirable.

load_tile.cuh:207:

        template <int LOAD, int dummy>
        struct Iterate<LOAD, LOAD_VEC_SIZE, dummy>

The error I'm getting is:

../../gunrock/util/io/load_tile.cuh:207:23: error: non-type template argument
      depends on a template parameter of the partial specialization

So:

http://stackoverflow.com/questions/8866881/nested-template-specialization-depending-on-enclosing-template-parameters

What's a little puzzling to me is that LOAD, LOAD_VEC_SIZE, and dummy
are all "simple identifier"s, as best I can tell.

@yzhwang notes "Because LOAD_VEC_SIZE is computed from:
LOAD_VEC_SIZE = 1 << LOG_LOAD_VEC_SIZE
and LOG_LOAD_VEC_SIZE is the argument of struct LoadTile,
which is the parent struct of struct Iterate. I don't know if @dumerrill
will know this better because his b40c only runs on linux and I'm not
sure if he has tested it using clang."

Illegal memory access for BC computation on as-Skitter

Hi all,

I'm trying to run some Betweenness Centrality tests using the following graph: http://www.cise.ufl.edu/research/sparse/matrices/SNAP/as-Skitter.html

Here's how I'm running the program:

$ ./bin/test_bc_7.0_x86_64 market ~/github/graph-utils/graphs/as-Skitter.mtx --device=0 --src=-1 --quick

and here's the output up until the first error:

Reading directly from previously stored CSR arrays ...
Done reading (3s).

Degree Histogram (1696415 vertices, 22190596 edges):
Degree 0: 0 (0.00%)
Degree 2^0: 218399 (12.87%)
Degree 2^1: 447171 (26.36%)
Degree 2^2: 456022 (26.88%)
Degree 2^3: 318795 (18.79%)
Degree 2^4: 149829 (8.83%)
Degree 2^5: 68987 (4.07%)
Degree 2^6: 21819 (1.29%)
Degree 2^7: 9364 (0.55%)
Degree 2^8: 3714 (0.22%)
Degree 2^9: 1396 (0.08%)
Degree 2^10: 536 (0.03%)
Degree 2^11: 211 (0.01%)
Degree 2^12: 85 (0.01%)
Degree 2^13: 60 (0.00%)
Degree 2^14: 24 (0.00%)
Degree 2^15: 3 (0.00%)

[../../gunrock/app/bc/bc_enactor.cuh, 299] edge_map_forward::Kernel failed (CUDA error 77: an illegal memory access was encountered)

I think the problem comes from a specific source vertex, because the computation will work if I set src=5, for example. For what it's worth, I ran this using a Tesla K40c GPU and the CUDA 7.0 toolkit on Ubuntu 14.04.

If you need any more info from my end I'm happy to provide it.

Thanks!

-Adam

clang issue with the Enact template (simple_example)

[ 71%] Building NVCC (Device) object simple_example/CMakeFiles/simple_example.dir//./simple_example_generated_simple_example.cu.o
/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu:417:88: error:
      'Enact' following the 'template' keyword does not refer to a template
  ...INSTRUMENT> ::template Enact< gunrock::app::cc::CCProblem< VertexId, Siz...
                   ~~~~~~~~ ^~~~~
/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu:417:154: error:
      expected '(' for function-style cast or type construction
  ...Enact< gunrock::app::cc::CCProblem< VertexId, SizeT, Value, true> > (cc_...
            ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ^
/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu:535:91: error:
      'Enact' following the 'template' keyword does not refer to a template
  ...INSTRUMENT> ::template Enact< gunrock::app::bfs::BFSProblem< VertexId, S...
                   ~~~~~~~~ ^~~~~
/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu:535:173: error:
      expected '(' for function-style cast or type construction
  ...gunrock::app::bfs::BFSProblem< VertexId, SizeT, Value, true, false, false> >...
     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ^
/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu:626:88: error:
      'Enact' following the 'template' keyword does not refer to a template
  ...INSTRUMENT> ::template Enact< gunrock::app::bc::BCProblem< VertexId, Siz...
                   ~~~~~~~~ ^~~~~
/Users/jowens/Documents/working/gunrock/simple_example/simple_example.cu:626:161: error:
      expected '(' for function-style cast or type construction
  ...gunrock::app::bc::BCProblem< VertexId, SizeT, Value, true, false> > (con...
     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ^

OS X current build failures

These are all Boost-related errors (but that doesn't mean it's Boost's fault). The Gunrock library builds after inlining all header functions. Some tests build, some do not. Summary below.

CMake Error at betweenness_centrality_generated_test_bc.cu.o.cmake:264 (message):
  Error generating file
  /Users/jowens/Documents/working/gunrock-build/tests/bc/CMakeFiles/betweenness_centrality.dir//./betweenness_centrality_generated_test_bc.cu.o


make[2]: *** [tests/bc/CMakeFiles/betweenness_centrality.dir/betweenness_centrality_generated_test_bc.cu.o] Error 1
make[2]: Target `tests/bc/CMakeFiles/betweenness_centrality.dir/depend' not remade because of errors.
make[1]: *** [tests/bc/CMakeFiles/betweenness_centrality.dir/all] Error 2
CMake Error at connected_component_generated_test_cc.cu.o.cmake:264 (message):
  Error generating file
  /Users/jowens/Documents/working/gunrock-build/tests/cc/CMakeFiles/connected_component.dir//./connected_component_generated_test_cc.cu.o


make[2]: *** [tests/cc/CMakeFiles/connected_component.dir/connected_component_generated_test_cc.cu.o] Error 1
make[2]: Target `tests/cc/CMakeFiles/connected_component.dir/depend' not remade because of errors.
make[1]: *** [tests/cc/CMakeFiles/connected_component.dir/all] Error 2
CMake Error at page_rank_generated_test_pr.cu.o.cmake:264 (message):
  Error generating file
  /Users/jowens/Documents/working/gunrock-build/tests/pr/CMakeFiles/page_rank.dir//./page_rank_generated_test_pr.cu.o


make[2]: *** [tests/pr/CMakeFiles/page_rank.dir/page_rank_generated_test_pr.cu.o] Error 1
make[2]: Target `tests/pr/CMakeFiles/page_rank.dir/depend' not remade because of errors.
make[1]: *** [tests/pr/CMakeFiles/page_rank.dir/all] Error 2
CMake Error at single_source_shortest_path_generated_test_sssp.cu.o.cmake:264 (message):
  Error generating file
  /Users/jowens/Documents/working/gunrock-build/tests/sssp/CMakeFiles/single_source_shortest_path.dir//./single_source_shortest_path_generated_test_sssp.cu.o


make[2]: *** [tests/sssp/CMakeFiles/single_source_shortest_path.dir/single_source_shortest_path_generated_test_sssp.cu.o] Error 1
make[2]: Target `tests/sssp/CMakeFiles/single_source_shortest_path.dir/depend' not remade because of errors.
make[1]: *** [tests/sssp/CMakeFiles/single_source_shortest_path.dir/all] Error 2
CMake Error at minimum_spanning_tree_generated_test_mst.cu.o.cmake:264 (message):
  Error generating file
  /Users/jowens/Documents/working/gunrock-build/tests/mst/CMakeFiles/minimum_spanning_tree.dir//./minimum_spanning_tree_generated_test_mst.cu.o


make[2]: *** [tests/mst/CMakeFiles/minimum_spanning_tree.dir/minimum_spanning_tree_generated_test_mst.cu.o] Error 1
make[2]: Target `tests/mst/CMakeFiles/minimum_spanning_tree.dir/depend' not remade because of errors.
make[1]: *** [tests/mst/CMakeFiles/minimum_spanning_tree.dir/all] Error 2
make[1]: Target `all' not remade because of errors.
make: *** [all] Error 2
make: Target `default_target' not remade because of errors.

TODOs before 0.3 release.

Primitive interface (MST, SALSA, WTF, HITS)
Update documentations, add release notes, change-log and some other docs.
Update all the run.sh, ppopp-test.sh under test/primitive dir.
Git remove broken/template primitives in master branch (MIS, VIS and Template).

~~- [ ] Hide memory info printfs in 5 primitives. (optional?)~~

Link why-gunrock.md into the top-level document.

bc algorithm

Hi,
I am trying to perform your bc implementation.
/bin/test_bc_6.5_x86_64 market ../../dataset/small/test_bc.mtx

I have several errors:
Computing reference value ...
CPU BC finished in 0.000000 msec.
[../../gunrock/app/bc/bc_enactor.cuh, 299] edge_map_forward::Kernel failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_enactor.cuh, 512] filter_forward::Kernel failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 437] BC Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_problem.cuh, 429] BCProblem cudaMemcpy data_slices to d_data_slices failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 436] BC Problem Data Reset Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_enactor.cuh, 109] BCEnactor cudaBindTexture row_offset_tex_ref failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 437] BC Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_problem.cuh, 429] BCProblem cudaMemcpy data_slices to d_data_slices failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 436] BC Problem Data Reset Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_enactor.cuh, 109] BCEnactor cudaBindTexture row_offset_tex_ref failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 437] BC Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_problem.cuh, 429] BCProblem cudaMemcpy data_slices to d_data_slices failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 436] BC Problem Data Reset Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_enactor.cuh, 109] BCEnactor cudaBindTexture row_offset_tex_ref failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 437] BC Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_problem.cuh, 429] BCProblem cudaMemcpy data_slices to d_data_slices failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 436] BC Problem Data Reset Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_enactor.cuh, 109] BCEnactor cudaBindTexture row_offset_tex_ref failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 437] BC Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_problem.cuh, 429] BCProblem cudaMemcpy data_slices to d_data_slices failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 436] BC Problem Data Reset Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_enactor.cuh, 109] BCEnactor cudaBindTexture row_offset_tex_ref failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 437] BC Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_problem.cuh, 429] BCProblem cudaMemcpy data_slices to d_data_slices failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 436] BC Problem Data Reset Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_enactor.cuh, 109] BCEnactor cudaBindTexture row_offset_tex_ref failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 437] BC Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/bc/bc_problem.cuh, 168] BCProblem cudaMemcpy d_bc_values failed (CUDA error 77: an illegal memory access was encountered)
[test_bc.cu, 452] BC Problem Data Extraction Failed (CUDA error 77: an illegal memory access was encountered)

Does your bc implementation work ?
Thanks,

Flavio

[--iteration-num=<num>] flag doesn't work.

Setting --iteration-num= flag (num > 1) will cause testing never end (hang in the 2nd test run). This call cannot stop: util::GRError(enactor->Enact(src, traversal_mode).

Cmake error configuring v0.3

I copied the cmake output here.

Need raw data for comparison against other engines

I've now recreated several graphs using vega-lite and Gunrock data (input is Gunrock json files). This is good. I now would like to start using non-Gunrock data. I am not finding the raw data from which we
made the tables in the current version of the Gunrock paper, except for a "wee-table.xlsx", which isn't exactly what I want anyway.

Specifically, I'd like to be able to get runtime data (elapsed time, possibly MTEPS) for a given combination of {engine, algorithm, dataset}. Where do the scripts that generated runtime numbers live?

Seems likely I'll either rewrite the .sh scripts that generate them as python scripts that output JSON, or write a short python script that inputs the output of the sh script and outputs JSON.

Also, if these engines are on mario or luigi, I'll need to be able to get to their binaries, so hopefully those are runnable as me.

broken cmake build: missing simple_example mtx file

CMake Error at simple_example/CMakeLists.txt:9 (file):
  file COPY cannot find
  "/Users/jowens/Documents/working/gunrock/simple_example/bips98_606.mtx_ud".


-- Configuring incomplete, errors occurred!

Possible Bug with Direction-Optimized BFS

Hey guys, I was running some tests with direction-optimized BFS and I noticed a problem with kron scale20 from source nodes 1 and 3. Strangely enough, the other nodes seem to be fine.

./bin/test_dobfs_7.0_x86_64 market /data/gunrock_dataset/large/kron_g500-logn20/kron_g500-logn20.mtx --src=0 --idempotence
./bin/test_dobfs_7.0_x86_64 market /data/gunrock_dataset/large/kron_g500-logn20/kron_g500-logn20.mtx --src=1 --idempotence
[../../gunrock/app/dobfs/dobfs_enactor.cuh, 525] BFSEnactor cudaEventSynchronize throttle_event failed (CUDA error 77: an illegal memory access was encountered)
[test_dobfs.cu, 399] DOBFS Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 186] DOBFSProblem cudaMemcpy d_labels failed (CUDA error 77: an illegal memory access was encountered)
[test_dobfs.cu, 410] DOBFS Problem Data Extraction Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 145] GpuSlice cudaFree d_labels failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 146] GpuSlice cudaFree d_preds failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 147] GpuSlice cudaFree d_frontier_map_in failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 148] GpuSlice cudaFree d_frontier_map_out failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 149] GpuSlice cudaFree d_index_queue failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 150] GpuSlice cudaFree d_index_queue failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 151] GpuSlice cudaFree data_slices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 126] GpuSlice cudaFree d_row_offsets failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 127] GpuSlice cudaFree d_column_indices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 128] GpuSlice cudaFree d_column_offsets failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 129] GpuSlice cudaFree d_row_indices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_enactor.cuh, 150] DOBFSEnactor cudaFreeHost done failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_enactor.cuh, 153] DOBFSEnactor cudaEventDestroy throttle_event failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 120] EnactorBase cudaFree d_node_locks failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 121] EnactorBase cudaFree d_node_locks_out failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed: (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed: (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/cta_work_progress.cuh, 230] CtaWorkProgress cudaFree d_counters failed: (CUDA error 77: an illegal memory access was encountered)
kron_g500-logn20
./bin/test_dobfs_7.0_x86_64 market /data/gunrock_dataset/large/kron_g500-logn20/kron_g500-logn20.mtx --src=2 --idempotence
kron_g500-logn20
./bin/test_dobfs_7.0_x86_64 market /data/gunrock_dataset/large/kron_g500-logn20/kron_g500-logn20.mtx --src=3 --idempotence
[../../gunrock/app/dobfs/dobfs_enactor.cuh, 525] BFSEnactor cudaEventSynchronize throttle_event failed (CUDA error 77: an illegal memory access was encountered)
[test_dobfs.cu, 399] DOBFS Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 186] DOBFSProblem cudaMemcpy d_labels failed (CUDA error 77: an illegal memory access was encountered)
[test_dobfs.cu, 410] DOBFS Problem Data Extraction Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 145] GpuSlice cudaFree d_labels failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 146] GpuSlice cudaFree d_preds failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 147] GpuSlice cudaFree d_frontier_map_in failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 148] GpuSlice cudaFree d_frontier_map_out failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 149] GpuSlice cudaFree d_index_queue failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 150] GpuSlice cudaFree d_index_queue failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_problem.cuh, 151] GpuSlice cudaFree data_slices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 126] GpuSlice cudaFree d_row_offsets failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 127] GpuSlice cudaFree d_column_indices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 128] GpuSlice cudaFree d_column_offsets failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 129] GpuSlice cudaFree d_row_indices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_enactor.cuh, 150] DOBFSEnactor cudaFreeHost done failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/dobfs/dobfs_enactor.cuh, 153] DOBFSEnactor cudaEventDestroy throttle_event failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 120] EnactorBase cudaFree d_node_locks failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 121] EnactorBase cudaFree d_node_locks_out failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed: (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed: (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/cta_work_progress.cuh, 230] CtaWorkProgress cudaFree d_counters failed: (CUDA error 77: an illegal memory access was encountered)
kron_g500-logn20
./bin/test_dobfs_7.0_x86_64 market /data/gunrock_dataset/large/kron_g500-logn20/kron_g500-logn20.mtx --src=4 --idempotence
kron_g500-logn20
./bin/test_dobfs_7.0_x86_64 market /data/gunrock_dataset/large/kron_g500-logn20/kron_g500-logn20.mtx --src=5 --idempotence
kron_g500-logn20
./bin/test_dobfs_7.0_x86_64 market /data/gunrock_dataset/large/kron_g500-logn20/kron_g500-logn20.mtx --src=6 --idempotence

Regards,
Carl

Multi-GPU Pagerank algorithm is returning incorrect results

Calling the Mgpu-pr algorithm with --device=0,0 or --device=0,1 results in the output being inconsistent with the CPU reference code.

Hang in BC, observed only on mgpu dev/merge branch

using graph: test_bc.mx when setting --src=-1 no matter whether add undirected or setting device.

[yuduo@mario]bc$ bin/test_bc_7.0_x86_64 market /data/gunrock_dataset/small/test_bc.mtx --src=-1 --device=0,1
Using 2 gpus:  0  1
src = -1
partition_method = random
Partition begin. seed=1436459662
1            out_offsets = 0, 4, 7  
0            out_offsets = 0, 3, 7
0            out_counter = 3, 4, 7
0            in_counter  = 3, 3, 6
1            out_counter = 4, 3, 7
1            in_counter  = 4, 4, 8
partition end. (0.350617 ms)
Computing reference value ...
CPU BC finished in 0.017811 msec.
0            in_counter = 3, 3, 6
1            in_counter = 4, 4, 8
__________________________
0            in_counter = 3, 3, 6
1            in_counter = 4, 4, 8
Iteration entered
Iteration entered
1    0   0   queue3      oversize :  6 ->    9
1    1   0   queue3      oversize :  6 ->    9
0    1   0   queue3      oversize :  5 ->    11 
0    1   0   queue3      oversize :  5 ->    9
1    2   0   queue3      oversize :  9 ->    10
Iteration entered
Iteration entered
0            in_counter = 3, 3, 6
1            in_counter = 4, 4, 8

and hang here. didn't try to add other flag to debug through.

master branch works fine:

[yuduo@mario]bc$ bin/test_bc_7.0_x86_64 market /data/gunrock_dataset/small/test_bc.mtx --src=-1
Computing reference value ...
CPU BC finished in 0.000000 msec.
Validity BC Value: 
CORRECT
[0:1.000000,0.500000 1:1.000000,0.500000 2:3.000000,2.333333 3:1.000000,1.833333        4:1.000000,1.833333 5:1.000000,0.666667 6:1.000000,0.333333 ]
GPU BC finished in 4.603424 msec.

simple_example fails compilation (boost 1.55 on OS X)

Not a high priority. Just checking status on OS X.

mission-burrito 4139$ pwd
/Users/jowens/Documents/working/gunrock/simple_example
mission-burrito 4093$ make
mkdir -p bin
"/Developer/NVIDIA/CUDA-5.5/bin/nvcc"  -gencode=arch=compute_30,code=\"sm_30,compute_30\" -o bin/simple_example_5.5_x86_64 simple_example.cu -Xptxas -v -Xcudafe -# -Xptxas -abi=no -m64 -I"/Developer/NVIDIA/CUDA-5.5/bin/../include" -I"/opt/local/include" -I.. -I../.. -lcuda -O3
/opt/local/include/boost/type_traits/is_abstract.hpp(72): error: identifier "__is_abstract" is undefined

/opt/local/include/boost/type_traits/is_abstract.hpp(72): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/is_abstract.hpp(72): error: type name is not allowed

/opt/local/include/boost/type_traits/is_convertible.hpp(484): error: identifier "__is_convertible_to" is undefined

/opt/local/include/boost/type_traits/is_convertible.hpp(484): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/is_convertible.hpp(484): error: type name is not allowed

/opt/local/include/boost/type_traits/is_convertible.hpp(484): error: type name is not allowed

/opt/local/include/boost/type_traits/is_class.hpp(123): error: identifier "__is_class" is undefined

/opt/local/include/boost/type_traits/is_class.hpp(123): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/is_class.hpp(123): error: type name is not allowed

/opt/local/include/boost/type_traits/is_enum.hpp(181): error: identifier "__is_enum" is undefined

/opt/local/include/boost/type_traits/is_enum.hpp(181): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/is_enum.hpp(181): error: type name is not allowed

/opt/local/include/boost/type_traits/has_trivial_destructor.hpp(28): error: identifier "__has_trivial_destructor" is undefined

/opt/local/include/boost/type_traits/has_trivial_destructor.hpp(28): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_trivial_destructor.hpp(28): error: type name is not allowed

/opt/local/include/boost/type_traits/has_trivial_copy.hpp(31): error: identifier "__has_trivial_copy" is undefined

/opt/local/include/boost/type_traits/has_trivial_copy.hpp(31): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_trivial_copy.hpp(31): error: type name is not allowed

/opt/local/include/boost/type_traits/has_nothrow_copy.hpp(24): error: identifier "__has_nothrow_copy" is undefined

/opt/local/include/boost/type_traits/has_nothrow_copy.hpp(24): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_nothrow_copy.hpp(24): error: type name is not allowed

/opt/local/include/boost/type_traits/has_trivial_assign.hpp(32): error: identifier "__has_trivial_assign" is undefined

/opt/local/include/boost/type_traits/has_trivial_assign.hpp(32): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_trivial_assign.hpp(32): error: type name is not allowed

/opt/local/include/boost/type_traits/has_nothrow_assign.hpp(26): error: identifier "__has_nothrow_assign" is undefined

/opt/local/include/boost/type_traits/has_nothrow_assign.hpp(26): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_nothrow_assign.hpp(26): error: type name is not allowed

/opt/local/include/boost/assert.hpp(102): warning: unknown attribute "__attribute__"

/opt/local/include/boost/functional/hash/hash.hpp(421): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (__nv_bool)

/opt/local/include/boost/functional/hash/hash.hpp(422): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (char)

/opt/local/include/boost/functional/hash/hash.hpp(423): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (signed char)

/opt/local/include/boost/functional/hash/hash.hpp(424): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (unsigned char)

/opt/local/include/boost/functional/hash/hash.hpp(426): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (wchar_t)

/opt/local/include/boost/functional/hash/hash.hpp(428): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (short)

/opt/local/include/boost/functional/hash/hash.hpp(429): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (unsigned short)

/opt/local/include/boost/functional/hash/hash.hpp(430): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (int)

/opt/local/include/boost/functional/hash/hash.hpp(431): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (unsigned int)

/opt/local/include/boost/functional/hash/hash.hpp(432): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (long)

/opt/local/include/boost/functional/hash/hash.hpp(433): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::basic_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (unsigned long)

/opt/local/include/boost/functional/hash/hash.hpp(435): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            function template "boost::hash_detail::float_numbers<T>::type boost::hash_value(T)"
            argument types are: (float)

/opt/local/include/boost/functional/hash/hash.hpp(436): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            function template "boost::hash_detail::float_numbers<T>::type boost::hash_value(T)"
            argument types are: (double)

/opt/local/include/boost/functional/hash/hash.hpp(437): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            function template "boost::hash_detail::float_numbers<T>::type boost::hash_value(T)"
            argument types are: (long double)

/opt/local/include/boost/functional/hash/hash.hpp(445): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::long_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (boost::long_long_type)

/opt/local/include/boost/functional/hash/hash.hpp(446): error: more than one instance of overloaded function "boost::hash_value" matches the argument list:
            function template "boost::hash_detail::ulong_numbers<T>::type boost::hash_value(T)"
            function template "boost::enable_if<boost::is_enum<T>, size_t>::type boost::hash_value(T)"
            argument types are: (boost::ulong_long_type)

/opt/local/include/boost/type_traits/has_trivial_constructor.hpp(28): error: identifier "__has_trivial_constructor" is undefined

/opt/local/include/boost/type_traits/has_trivial_constructor.hpp(28): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_trivial_constructor.hpp(28): error: type name is not allowed

/opt/local/include/boost/type_traits/has_nothrow_constructor.hpp(24): error: identifier "__has_nothrow_constructor" is undefined

/opt/local/include/boost/type_traits/has_nothrow_constructor.hpp(24): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_nothrow_constructor.hpp(24): error: type name is not allowed

/opt/local/include/boost/type_traits/is_base_and_derived.hpp(228): error: identifier "__is_base_of" is undefined

/opt/local/include/boost/type_traits/is_base_and_derived.hpp(228): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/is_base_and_derived.hpp(228): error: type name is not allowed

/opt/local/include/boost/type_traits/is_base_and_derived.hpp(228): error: type name is not allowed

/opt/local/include/boost/type_traits/has_virtual_destructor.hpp(20): error: identifier "__has_virtual_destructor" is undefined

/opt/local/include/boost/type_traits/has_virtual_destructor.hpp(20): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/has_virtual_destructor.hpp(20): error: type name is not allowed

/opt/local/include/boost/type_traits/is_polymorphic.hpp(106): error: identifier "__is_polymorphic" is undefined

/opt/local/include/boost/type_traits/is_polymorphic.hpp(106): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/is_polymorphic.hpp(106): error: type name is not allowed

/opt/local/include/boost/type_traits/is_union.hpp(43): error: identifier "__is_union" is undefined

/opt/local/include/boost/type_traits/is_union.hpp(43): error: function call is not allowed in a constant expression

/opt/local/include/boost/type_traits/is_union.hpp(43): error: type name is not allowed

/opt/local/include/boost/smart_ptr/detail/sp_counted_base_gcc_x86.hpp(49): warning: "cc" clobber ignored

/opt/local/include/boost/smart_ptr/detail/sp_counted_base_gcc_x86.hpp(65): warning: "cc" clobber ignored

/opt/local/include/boost/smart_ptr/detail/sp_counted_base_gcc_x86.hpp(91): warning: "cc" clobber ignored

/opt/local/include/boost/smart_ptr/detail/sp_counted_base_gcc_x86.hpp(75): warning: variable "tmp" was set but never used

simple_example.cu(461): error: template instantiation resulted in unexpected function type of "int (const boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS> &, int *, boost::graph::detail::no_parameter)" (the meaning of a name may have changed since the template declaration -- the type of the template is "boost::property_traits<ComponentMap>::value_type (const Graph &, ComponentMap, boost::enable_if_c<boost::is_base_and_derived<boost::vertex_list_graph_tag, boost::graph_traits<Graph>::traversal_category>::value, boost::graph::detail::no_parameter>::type)")
          detected during:
            instantiation of "boost::connected_components" based on template arguments <boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, int *>
(461): here
            instantiation of "void RunTests<VertexId,Value,SizeT,INSTRUMENT>(const gunrock::Csr<VertexId, Value, SizeT> &, int, int, double) [with VertexId=int, Value=float, SizeT=int, INSTRUMENT=true]"
(816): here
            instantiation of "void RunTests(gunrock::Csr<VertexId, Value, SizeT> &, gunrock::util::CommandLineArgs &) [with VertexId=int, Value=float, SizeT=int]"
(884): here

/opt/local/include/boost/iterator/iterator_facade.hpp(842): error: no instance of overloaded function "boost::iterator_core_access::distance_from" matches the argument list
            argument types are: (const boost::range_detail::integer_iterator<size_t>, const boost::range_detail::integer_iterator<size_t>, boost::is_convertible<boost::range_detail::integer_iterator<size_t>, boost::range_detail::integer_iterator<size_t>>)
          detected during:
            instantiation of "boost::detail::enable_if_interoperable<Derived1, Derived2, boost::mpl::apply2<boost::detail::choose_difference_type, Derived1, Derived2>::type>::type boost::operator-(const boost::iterator_facade<Derived1, V1, TC1, Reference1, Difference1> &, const boost::iterator_facade<Derived2, V2, TC2, Reference2, Difference2> &) [with Derived1=boost::range_detail::integer_iterator<size_t>, V1=size_t, TC1=boost::random_access_traversal_tag, Reference1=size_t, Difference1=ptrdiff_t, Derived2=boost::range_detail::integer_iterator<size_t>, V2=size_t, TC2=boost::random_access_traversal_tag, Reference2=size_t, Difference2=ptrdiff_t]"
/opt/local/include/boost/range/iterator_range_core.hpp(255): here
            instantiation of "boost::iterator_range<IteratorT>::difference_type boost::iterator_range<IteratorT>::size() const [with IteratorT=boost::range_detail::integer_iterator<size_t>]"
/opt/local/include/boost/graph/detail/adjacency_list.hpp(1568): here
            instantiation of "Config::vertices_size_type boost::num_vertices(const boost::adj_list_helper<Config, Base> &) [with Config=boost::detail::adj_list_gen<boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>::config, Base=boost::undirected_graph_helper<boost::detail::adj_list_gen<boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>::config>]"
/opt/local/include/boost/graph/detail/adjacency_list.hpp(2179): here
            instantiation of "std::pair<Config::edge_descriptor, __nv_bool> boost::add_edge(Config::vertex_descriptor, Config::vertex_descriptor, const Config::edge_property_type &, boost::vec_adj_list_impl<Graph, Config, Base> &) [with Graph=boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, Config=boost::detail::adj_list_gen<boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>::config, Base=boost::undirected_graph_helper<boost::detail::adj_list_gen<boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>::config>]"
/opt/local/include/boost/graph/detail/adjacency_list.hpp(2191): here
            instantiation of "std::pair<Config::edge_descriptor, __nv_bool> boost::add_edge(Config::vertex_descriptor, Config::vertex_descriptor, boost::vec_adj_list_impl<Graph, Config, Base> &) [with Graph=boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, Config=boost::detail::adj_list_gen<boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>::config, Base=boost::undirected_graph_helper<boost::detail::adj_list_gen<boost::adjacency_list<boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>, boost::vecS, boost::vecS, boost::undirectedS, boost::no_property, boost::no_property, boost::no_property, boost::listS>::config>]"
simple_example.cu(456): here
            instantiation of "unsigned int RefCPUCC(SizeT *, VertexId *, int, int *) [with VertexId=int, SizeT=int]"
simple_example.cu(528): here
            instantiation of "void RunTests<VertexId,Value,SizeT,INSTRUMENT>(const gunrock::Csr<VertexId, Value, SizeT> &, int, int, double) [with VertexId=int, Value=float, SizeT=int, INSTRUMENT=true]"
simple_example.cu(816): here
            instantiation of "void RunTests(gunrock::Csr<VertexId, Value, SizeT> &, gunrock::util::CommandLineArgs &) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(884): here

/opt/local/include/boost/iterator/iterator_facade.hpp(833): error: no instance of overloaded function "boost::iterator_core_access::equal" matches the argument list
            argument types are: (const boost::range_detail::integer_iterator<size_t>, const boost::range_detail::integer_iterator<size_t>, boost::is_convertible<boost::range_detail::integer_iterator<size_t>, boost::range_detail::integer_iterator<size_t>>)
          detected during:
            instantiation of "boost::detail::enable_if_interoperable<Derived1, Derived2, boost::mpl::apply2<boost::detail::always_bool2, Derived1, Derived2>::type>::type boost::operator!=(const boost::iterator_facade<Derived1, V1, TC1, Reference1, Difference1> &, const boost::iterator_facade<Derived2, V2, TC2, Reference2, Difference2> &) [with Derived1=boost::range_detail::integer_iterator<size_t>, V1=size_t, TC1=boost::random_access_traversal_tag, Reference1=size_t, Difference1=ptrdiff_t, Derived2=boost::range_detail::integer_iterator<size_t>, V2=size_t, TC2=boost::random_access_traversal_tag, Reference2=size_t, Difference2=ptrdiff_t]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(249): here
            instantiation of "void boost::detail::graph::init_centrality_map(std::pair<Iter, Iter>, Centrality) [with Iter=boost::range_detail::integer_iterator<size_t>, Centrality=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(302): here
            instantiation of "void boost::detail::graph::brandes_betweenness_centrality_impl(const Graph &, CentralityMap, EdgeCentralityMap, IncomingMap, DistanceMap, DependencyMap, PathCountMap, VertexIndexMap, ShortestPaths) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, IncomingMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> *, std::vector<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::allocator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> &>, DistanceMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, DependencyMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, PathCountMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<size_t *, std::vector<size_t, std::allocator<size_t>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, unsigned long, unsigned long &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, ShortestPaths=boost::detail::graph::brandes_unweighted_shortest_paths]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(382): here
            instantiation of "void boost::brandes_betweenness_centrality(const Graph &, CentralityMap, EdgeCentralityMap, IncomingMap, DistanceMap, DependencyMap, PathCountMap, VertexIndexMap, boost::enable_if_c<boost::is_base_and_derived<boost::vertex_list_graph_tag, boost::graph_traits<Graph>::traversal_category>::value, boost::graph::detail::no_parameter>::type) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, IncomingMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> *, std::vector<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::allocator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> &>, DistanceMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, DependencyMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, PathCountMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<size_t *, std::vector<size_t, std::allocator<size_t>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, unsigned long, unsigned long &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(479): here
            instantiation of "void boost::detail::graph::brandes_betweenness_centrality_dispatch2(const Graph &, CentralityMap, EdgeCentralityMap, VertexIndexMap) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(562): here
            instantiation of "void boost::brandes_betweenness_centrality(const Graph &, CentralityMap, EdgeCentralityMap, boost::enable_if_c<boost::is_base_and_derived<boost::vertex_list_graph_tag, boost::graph_traits<Graph>::traversal_category>::value, boost::graph::detail::no_parameter>::type) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>]"
simple_example.cu(421): here
            instantiation of "void RefCPUBC(const gunrock::Csr<VertexId, Value, SizeT> &, Value *, VertexId) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(728): here
            instantiation of "void RunTests<VertexId,Value,SizeT,INSTRUMENT>(const gunrock::Csr<VertexId, Value, SizeT> &, int, int, double) [with VertexId=int, Value=float, SizeT=int, INSTRUMENT=true]"
simple_example.cu(816): here
            instantiation of "void RunTests(gunrock::Csr<VertexId, Value, SizeT> &, gunrock::util::CommandLineArgs &) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(884): here

/opt/local/include/boost/iterator/iterator_facade.hpp(833): error: no instance of overloaded function "boost::iterator_core_access::equal" matches the argument list
            argument types are: (const boost::detail::undirected_edge_iter<std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, const boost::detail::undirected_edge_iter<std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, boost::is_convertible<boost::detail::undirected_edge_iter<std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, boost::detail::undirected_edge_iter<std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>>)
          detected during:
            instantiation of "boost::detail::enable_if_interoperable<Derived1, Derived2, boost::mpl::apply2<boost::detail::always_bool2, Derived1, Derived2>::type>::type boost::operator!=(const boost::iterator_facade<Derived1, V1, TC1, Reference1, Difference1> &, const boost::iterator_facade<Derived2, V2, TC2, Reference2, Difference2> &) [with Derived1=boost::detail::undirected_edge_iter<std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, V1=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, TC1=std::bidirectional_iterator_tag, Reference1=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, Difference1=ptrdiff_t, Derived2=boost::detail::undirected_edge_iter<std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, V2=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, TC2=std::bidirectional_iterator_tag, Reference2=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, Difference2=ptrdiff_t]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(249): here
            instantiation of "void boost::detail::graph::init_centrality_map(std::pair<Iter, Iter>, Centrality) [with Iter=boost::detail::undirected_edge_iter<std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, Centrality=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(303): here
            instantiation of "void boost::detail::graph::brandes_betweenness_centrality_impl(const Graph &, CentralityMap, EdgeCentralityMap, IncomingMap, DistanceMap, DependencyMap, PathCountMap, VertexIndexMap, ShortestPaths) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, IncomingMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> *, std::vector<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::allocator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> &>, DistanceMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, DependencyMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, PathCountMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<size_t *, std::vector<size_t, std::allocator<size_t>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, unsigned long, unsigned long &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, ShortestPaths=boost::detail::graph::brandes_unweighted_shortest_paths]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(382): here
            instantiation of "void boost::brandes_betweenness_centrality(const Graph &, CentralityMap, EdgeCentralityMap, IncomingMap, DistanceMap, DependencyMap, PathCountMap, VertexIndexMap, boost::enable_if_c<boost::is_base_and_derived<boost::vertex_list_graph_tag, boost::graph_traits<Graph>::traversal_category>::value, boost::graph::detail::no_parameter>::type) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, IncomingMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> *, std::vector<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::allocator<std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>>, std::vector<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, std::allocator<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>> &>, DistanceMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, DependencyMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, PathCountMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<size_t *, std::vector<size_t, std::allocator<size_t>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, unsigned long, unsigned long &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(479): here
            instantiation of "void boost::detail::graph::brandes_betweenness_centrality_dispatch2(const Graph &, CentralityMap, EdgeCentralityMap, VertexIndexMap) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(562): here
            instantiation of "void boost::brandes_betweenness_centrality(const Graph &, CentralityMap, EdgeCentralityMap, boost::enable_if_c<boost::is_base_and_derived<boost::vertex_list_graph_tag, boost::graph_traits<Graph>::traversal_category>::value, boost::graph::detail::no_parameter>::type) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>]"
simple_example.cu(421): here
            instantiation of "void RefCPUBC(const gunrock::Csr<VertexId, Value, SizeT> &, Value *, VertexId) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(728): here
            instantiation of "void RunTests<VertexId,Value,SizeT,INSTRUMENT>(const gunrock::Csr<VertexId, Value, SizeT> &, int, int, double) [with VertexId=int, Value=float, SizeT=int, INSTRUMENT=true]"
simple_example.cu(816): here
            instantiation of "void RunTests(gunrock::Csr<VertexId, Value, SizeT> &, gunrock::util::CommandLineArgs &) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(884): here

/opt/local/include/boost/iterator/iterator_facade.hpp(832): error: no instance of overloaded function "boost::iterator_core_access::equal" matches the argument list
            argument types are: (const boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, const boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, boost::is_convertible<boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>>)
          detected during:
            instantiation of "boost::detail::enable_if_interoperable<Derived1, Derived2, boost::mpl::apply2<boost::detail::always_bool2, Derived1, Derived2>::type>::type boost::operator==(const boost::iterator_facade<Derived1, V1, TC1, Reference1, Difference1> &, const boost::iterator_facade<Derived2, V2, TC2, Reference2, Difference2> &) [with Derived1=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, V1=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, TC1=std::bidirectional_iterator_tag, Reference1=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, Difference1=ptrdiff_t, Derived2=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, V2=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, TC2=std::bidirectional_iterator_tag, Reference2=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, Difference2=ptrdiff_t]"
/opt/local/include/boost/concept_check.hpp(237): here
            instantiation of "boost::EqualityComparable<TT>::~EqualityComparable() [with TT=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>]"
/opt/local/include/boost/concept_check.hpp(516): here
            instantiation of "boost::InputIterator<TT>::~InputIterator() [with TT=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>]"
/opt/local/include/boost/concept/detail/general.hpp(38): here
            instantiation of "void boost::concepts::requirement<boost::concepts::failed ************Model::************>::failed() [with Model=boost::InputIterator<boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>>]"
/opt/local/include/boost/graph/graph_concepts.hpp(54): here
            instantiation of "boost::concepts::MultiPassInputIterator<T>::~MultiPassInputIterator() [with T=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>]"
/opt/local/include/boost/concept/detail/general.hpp(38): here
            [ 10 instantiation contexts not shown ]
            instantiation of "void boost::detail::graph::brandes_betweenness_centrality_dispatch2(const Graph &, CentralityMap, EdgeCentralityMap, VertexIndexMap) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(562): here
            instantiation of "void boost::brandes_betweenness_centrality(const Graph &, CentralityMap, EdgeCentralityMap, boost::enable_if_c<boost::is_base_and_derived<boost::vertex_list_graph_tag, boost::graph_traits<Graph>::traversal_category>::value, boost::graph::detail::no_parameter>::type) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>]"
simple_example.cu(421): here
            instantiation of "void RefCPUBC(const gunrock::Csr<VertexId, Value, SizeT> &, Value *, VertexId) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(728): here
            instantiation of "void RunTests<VertexId,Value,SizeT,INSTRUMENT>(const gunrock::Csr<VertexId, Value, SizeT> &, int, int, double) [with VertexId=int, Value=float, SizeT=int, INSTRUMENT=true]"
simple_example.cu(816): here
            instantiation of "void RunTests(gunrock::Csr<VertexId, Value, SizeT> &, gunrock::util::CommandLineArgs &) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(884): here

/opt/local/include/boost/iterator/iterator_facade.hpp(833): error: no instance of overloaded function "boost::iterator_core_access::equal" matches the argument list
            argument types are: (const boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, const boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, boost::is_convertible<boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>>)
          detected during:
            instantiation of "boost::detail::enable_if_interoperable<Derived1, Derived2, boost::mpl::apply2<boost::detail::always_bool2, Derived1, Derived2>::type>::type boost::operator!=(const boost::iterator_facade<Derived1, V1, TC1, Reference1, Difference1> &, const boost::iterator_facade<Derived2, V2, TC2, Reference2, Difference2> &) [with Derived1=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, V1=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, TC1=std::bidirectional_iterator_tag, Reference1=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, Difference1=ptrdiff_t, Derived2=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>, V2=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, TC2=std::bidirectional_iterator_tag, Reference2=boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, Difference2=ptrdiff_t]"
/opt/local/include/boost/concept_check.hpp(238): here
            instantiation of "boost::EqualityComparable<TT>::~EqualityComparable() [with TT=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>]"
/opt/local/include/boost/concept_check.hpp(516): here
            instantiation of "boost::InputIterator<TT>::~InputIterator() [with TT=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>]"
/opt/local/include/boost/concept/detail/general.hpp(38): here
            instantiation of "void boost::concepts::requirement<boost::concepts::failed ************Model::************>::failed() [with Model=boost::InputIterator<boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>>]"
/opt/local/include/boost/graph/graph_concepts.hpp(54): here
            instantiation of "boost::concepts::MultiPassInputIterator<T>::~MultiPassInputIterator() [with T=boost::detail::out_edge_iter<std::_Rb_tree_const_iterator<boost::detail::stored_edge_iter<size_t, std::_List_iterator<boost::list_edge<size_t, EdgeProperties>>, EdgeProperties>>, size_t, boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, ptrdiff_t>]"
/opt/local/include/boost/concept/detail/general.hpp(38): here
            [ 10 instantiation contexts not shown ]
            instantiation of "void boost::detail::graph::brandes_betweenness_centrality_dispatch2(const Graph &, CentralityMap, EdgeCentralityMap, VertexIndexMap) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>, VertexIndexMap=boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>]"
/opt/local/include/boost/graph/betweenness_centrality.hpp(562): here
            instantiation of "void boost::brandes_betweenness_centrality(const Graph &, CentralityMap, EdgeCentralityMap, boost::enable_if_c<boost::is_base_and_derived<boost::vertex_list_graph_tag, boost::graph_traits<Graph>::traversal_category>::value, boost::graph::detail::no_parameter>::type) [with Graph=boost::adjacency_list<boost::setS, boost::vecS, boost::undirectedS, boost::no_property, EdgeProperties, boost::no_property, boost::listS>, CentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::vec_adj_list_vertex_id_map<boost::no_property, size_t>, double, double &>, EdgeCentralityMap=boost::iterator_property_map<__gnu_cxx::__normal_iterator<double *, std::vector<double, std::allocator<double>>>, boost::associative_property_map<std::map<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int, std::less<boost::detail::edge_desc_impl<boost::undirected_tag, size_t>>, std::allocator<std::pair<const boost::detail::edge_desc_impl<boost::undirected_tag, size_t>, int>>>>, double, double &>]"
simple_example.cu(421): here
            instantiation of "void RefCPUBC(const gunrock::Csr<VertexId, Value, SizeT> &, Value *, VertexId) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(728): here
            instantiation of "void RunTests<VertexId,Value,SizeT,INSTRUMENT>(const gunrock::Csr<VertexId, Value, SizeT> &, int, int, double) [with VertexId=int, Value=float, SizeT=int, INSTRUMENT=true]"
simple_example.cu(816): here
            instantiation of "void RunTests(gunrock::Csr<VertexId, Value, SizeT> &, gunrock::util::CommandLineArgs &) [with VertexId=int, Value=float, SizeT=int]"
simple_example.cu(884): here

Front end time                       1.83 (CPU)       2.00 (elapsed)
69 errors detected in the compilation of "/tmp/tmpxft_00012e3f_00000000-8_simple_example.cpp1.ii".
Total compilation time               1.88 (CPU)       2.00 (elapsed)
make: *** [bin/simple_example_5.5_x86_64] Error 2

SALSA primitive Advance V2E and E2V bug

SALSA uses V2E and E2V which causes memory access error on several datasets. Probably due to a series change of operator but only tested for V2V.

PageRank and Personal PageRank CPU reference algorithms

PageRank BGL reference algorithm fails to remove dangling nodes, causing inconsistent results with our GPU results. (Boost version 1.53.0). Also, need to find a personal PageRank reference algorithm (Currently no ref algorithm so the validity test will fail).

PageRank radixsort error on directed graphs.

[../../gunrock/app/pr/pr_enactor.cuh, 466] cub::DeviceRadixSort::SortPairsDescending failed (CUDA error 9: invalid configuration argument)
[test_pr.cu, 339] pr Problem Enact Failed (CUDA error 9: invalid configuration argument)

Add switch to disable unnecessary work in Filter

For operation we know we do not need to filter out anything, we should add a switch to allow users to disable scan+compact for some extra performance.

clang 3.5 and bfs example: template errors

test_bfs.cu:372:91: error: 'Enact' following the 'template' keyword does not
      refer to a template
  ...INSTRUMENT> ::template Enact< gunrock::app::bfs::BFSProblem< VertexId, S...
                   ~~~~~~~~ ^~~~~
test_bfs.cu:372:233: error: expected '(' for function-style cast or type
      construction
  ...gunrock::app::bfs::BFSProblem< VertexId, SizeT, Value, MARK_PREDECESSORS, ENABLE_IDEMPOTENCE, MARK_PREDECESSORS && ENABLE_IDEMPOTENCE> >...
     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ^

Minimum Spanning Tree got illegal memory access error on ak2010 dateset

@ydwu I'm testing MST today and found out this bug.

bin/test_mst_7.0_x86_64 market /data/gunrock_dataset/large/ak2010/ak2010.mtx | tee result
Reading from /data/gunrock_dataset/large/ak2010/ak2010.mtx:
  Parsing MARKET COO format (45292 nodes, 108549 directed edges)... Done parsing (0s).
  Converting 45292 vertices, 217098 directed edges (unordered tuples) to CSR format...
Done converting (0s).
[../../gunrock/app/mst/mst_enactor.cuh, 511] MSTProblem cudaMemcpy vertex_flag to d_vertex_flag failed (CUDA error 77: an illegal memory access was encountered)
[test_mst.cu, 253] MST Problem Enact Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 232] MSTProblem cudaMemcpy selector failed (CUDA error 77: an illegal memory access was encountered)
[test_mst.cu, 262] MST Problem Data Extraction Failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 143] GpuSlice cudaFree d_col_indices  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 146] GpuSlice cudaFree d_edge_weights failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 149] GpuSlice cudaFree d_reduced_vals failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 152] GpuSlice cudaFree  d_flags_array failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 155] GpuSlice cudaFree  d_keys_array  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 158] GpuSlice cudaFree d_temp_storage failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 161] GpuSlice cudaFree d_reduced_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 164] GpuSlice cudaFree  d_successors  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 167] GpuSlice cudaFree d_row_offsets  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 170] GpuSlice cudaFree d_vertex_flag  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 173] GpuSlice cudaFree d_origin_nodes failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 176] GpuSlice cudaFree d_supervtx_ids failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 179] GpuSlice cudaFree d_origin_edges failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 182] GpuSlice cudaFree  d_mst_output  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 185] GpuSlice cudaFree  d_edge_flags  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 188] GpuSlice cudaFree d_tmp_storage  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 191] GpuSlice cudaFree d_super_edges  failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_problem.cuh, 195] GpuSlice cudaFree data_slices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 126] GpuSlice cudaFree d_row_offsets failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 127] GpuSlice cudaFree d_column_indices failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 132] GpuSlice cudaFree frontier_queues.d_values failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 131] GpuSlice cudaFree frontier_queues.d_keys failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/problem_base.cuh, 132] GpuSlice cudaFree frontier_queues.d_values failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_enactor.cuh, 144] MSTEnactor cudaFreeHost done failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/mst/mst_enactor.cuh, 148] MSTEnactor cudaEventDestroy throttle_event failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 120] EnactorBase cudaFree d_node_locks failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/app/enactor_base.cuh, 121] EnactorBase cudaFree d_node_locks_out failed (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed:  (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/kernel_runtime_stats.cuh, 143] KernelRuntimeStatsLifetime cudaFree d_stat failed:  (CUDA error 77: an illegal memory access was encountered)
[../../gunrock/util/cta_work_progress.cuh, 230] CtaWorkProgress cudaFree d_counters failed:  (CUDA error 77: an illegal memory access was encountered)

MINIMUM SPANNING TREE TEST
GPU - Computation Complete in 0.000000 msec.

REFERENCE TEST
CPU - Computation Complete in 138.863480 msec.
INCORRECT. 
CPU Computed Total Weight = 5414746115
GPU Computed Total Weight = 50497260931659

Query

I am working on cuda implementation of breadth first search using gunrock.
Can you explain how the cuda threads are distributing their work among each other and the algorithmic details.
Also which are the relevant files to explore for getting to understand the 'gunrock bfs' implementation.

How can I convert file from .mtx format to .gr format ?
How can I convert file from .gr format to .mtx format ?
Can I get some performance numbers of performance of gunrock implementations on various GPU hardwares?

I would be grateful if you could help me in this,
Awaiting your reply,
Best regards,
-Mayank

CPU Reference (BGL) for Connected Component incorrect.

BGL Connected Component results inconsistent with Gunrock and NetworkX on generated twitter graphs.