Adding support for doing reduce on neighbor list

gunrock · Feb 4, 2015 · f63ac6f · f63ac6f · 1duo · Feb 4, 2015
1 parent 1d4aa2b
commit f63ac6f
Show file tree

Hide file tree

Showing 224 changed files with 106 additions and 13 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
diff --git a/FAQ.markdown b/FAQ.markdown
diff --git a/FAQ.md b/FAQ.md
diff --git a/LICENSE.TXT b/LICENSE.TXT
diff --git a/README.md b/README.md
diff --git a/cmake/FindBoostHeaders.cmake b/cmake/FindBoostHeaders.cmake
diff --git a/codesnaps/annotated_primitives/all.html b/codesnaps/annotated_primitives/all.html
diff --git a/codesnaps/annotated_primitives/annotated_primitives.html b/codesnaps/annotated_primitives/annotated_primitives.html
diff --git a/codesnaps/annotated_primitives/bfs_enactor.html b/codesnaps/annotated_primitives/bfs_enactor.html
diff --git a/codesnaps/annotated_primitives/bfs_functor.html b/codesnaps/annotated_primitives/bfs_functor.html
diff --git a/codesnaps/annotated_primitives/bfs_problem.html b/codesnaps/annotated_primitives/bfs_problem.html
diff --git a/codesnaps/annotated_primitives/pycco.css b/codesnaps/annotated_primitives/pycco.css
diff --git a/codesnaps/annotated_primitives/salsa_enactor.html b/codesnaps/annotated_primitives/salsa_enactor.html
diff --git a/codesnaps/annotated_primitives/salsa_functor.html b/codesnaps/annotated_primitives/salsa_functor.html
diff --git a/codesnaps/annotated_primitives/salsa_problem.html b/codesnaps/annotated_primitives/salsa_problem.html
diff --git a/codesnaps/bfs/bfs_enactor.cuh b/codesnaps/bfs/bfs_enactor.cuh
diff --git a/codesnaps/bfs/bfs_functor.cuh b/codesnaps/bfs/bfs_functor.cuh
diff --git a/codesnaps/bfs/bfs_problem.cuh b/codesnaps/bfs/bfs_problem.cuh
diff --git a/codesnaps/salsa/salsa_enactor.cuh b/codesnaps/salsa/salsa_enactor.cuh
diff --git a/codesnaps/salsa/salsa_functor.cuh b/codesnaps/salsa/salsa_functor.cuh
diff --git a/codesnaps/salsa/salsa_problem.cuh b/codesnaps/salsa/salsa_problem.cuh
diff --git a/dataset/large/Makefile b/dataset/large/Makefile
diff --git a/dataset/large/ak2010/Makefile b/dataset/large/ak2010/Makefile
diff --git a/dataset/large/belgium_osm/Makefile b/dataset/large/belgium_osm/Makefile
diff --git a/dataset/large/coAuthorsDBLP/Makefile b/dataset/large/coAuthorsDBLP/Makefile
diff --git a/dataset/large/common.mk b/dataset/large/common.mk
diff --git a/dataset/large/delaunay_n13/Makefile b/dataset/large/delaunay_n13/Makefile
diff --git a/dataset/large/delaunay_n21/Makefile b/dataset/large/delaunay_n21/Makefile
diff --git a/dataset/large/kron_g500-logn21/Makefile b/dataset/large/kron_g500-logn21/Makefile
diff --git a/dataset/large/soc-LiveJournal1/Makefile b/dataset/large/soc-LiveJournal1/Makefile
diff --git a/dataset/large/ufl.mk b/dataset/large/ufl.mk
diff --git a/dataset/large/webbase-1M/Makefile b/dataset/large/webbase-1M/Makefile
diff --git a/dataset/small/chesapeake.mtx b/dataset/small/chesapeake.mtx
diff --git a/dataset/small/test_bc.mtx b/dataset/small/test_bc.mtx
diff --git a/dataset/small/test_cc.mtx b/dataset/small/test_cc.mtx
diff --git a/dataset/small/test_pr.mtx b/dataset/small/test_pr.mtx
diff --git a/doc/building_gunrock.md b/doc/building_gunrock.md
diff --git a/doc/gunrock.doxygen b/doc/gunrock.doxygen
diff --git a/doc/programming_model.md b/doc/programming_model.md
diff --git a/doc/simple_example.dox b/doc/simple_example.dox
diff --git a/gunrock/CMakeLists.txt b/gunrock/CMakeLists.txt
diff --git a/gunrock/app/bc/bc_app.cu b/gunrock/app/bc/bc_app.cu
diff --git a/gunrock/app/bc/bc_enactor.cuh b/gunrock/app/bc/bc_enactor.cuh
diff --git a/gunrock/app/bc/bc_functor.cuh b/gunrock/app/bc/bc_functor.cuh
diff --git a/gunrock/app/bc/bc_problem.cuh b/gunrock/app/bc/bc_problem.cuh
diff --git a/gunrock/app/bfs/bfs_app.cu b/gunrock/app/bfs/bfs_app.cu
diff --git a/gunrock/app/bfs/bfs_enactor.cuh b/gunrock/app/bfs/bfs_enactor.cuh
diff --git a/gunrock/app/bfs/bfs_functor.cuh b/gunrock/app/bfs/bfs_functor.cuh
diff --git a/gunrock/app/bfs/bfs_problem.cuh b/gunrock/app/bfs/bfs_problem.cuh
diff --git a/gunrock/app/cc/cc_app.cu b/gunrock/app/cc/cc_app.cu
diff --git a/gunrock/app/cc/cc_enactor.cuh b/gunrock/app/cc/cc_enactor.cuh
diff --git a/gunrock/app/cc/cc_functor.cuh b/gunrock/app/cc/cc_functor.cuh
diff --git a/gunrock/app/cc/cc_problem.cuh b/gunrock/app/cc/cc_problem.cuh
diff --git a/gunrock/app/dobfs/dobfs_enactor.cuh b/gunrock/app/dobfs/dobfs_enactor.cuh
diff --git a/gunrock/app/dobfs/dobfs_functor.cuh b/gunrock/app/dobfs/dobfs_functor.cuh
diff --git a/gunrock/app/dobfs/dobfs_problem.cuh b/gunrock/app/dobfs/dobfs_problem.cuh
diff --git a/gunrock/app/enactor_base.cuh b/gunrock/app/enactor_base.cuh
diff --git a/gunrock/app/hits/hits_enactor.cuh b/gunrock/app/hits/hits_enactor.cuh
diff --git a/gunrock/app/hits/hits_functor.cuh b/gunrock/app/hits/hits_functor.cuh
diff --git a/gunrock/app/hits/hits_problem.cuh b/gunrock/app/hits/hits_problem.cuh
diff --git a/gunrock/app/mst/mst_enactor.cuh b/gunrock/app/mst/mst_enactor.cuh
diff --git a/gunrock/app/mst/mst_functor.cuh b/gunrock/app/mst/mst_functor.cuh
diff --git a/gunrock/app/mst/mst_problem.cuh b/gunrock/app/mst/mst_problem.cuh
diff --git a/gunrock/app/pr/pr_app.cu b/gunrock/app/pr/pr_app.cu
diff --git a/gunrock/app/pr/pr_enactor.cuh b/gunrock/app/pr/pr_enactor.cuh
diff --git a/gunrock/app/pr/pr_functor.cuh b/gunrock/app/pr/pr_functor.cuh
diff --git a/gunrock/app/pr/pr_problem.cuh b/gunrock/app/pr/pr_problem.cuh
diff --git a/gunrock/app/problem_base.cuh b/gunrock/app/problem_base.cuh
diff --git a/gunrock/app/salsa/salsa_enactor.cuh b/gunrock/app/salsa/salsa_enactor.cuh
diff --git a/gunrock/app/salsa/salsa_functor.cuh b/gunrock/app/salsa/salsa_functor.cuh
diff --git a/gunrock/app/salsa/salsa_problem.cuh b/gunrock/app/salsa/salsa_problem.cuh
diff --git a/gunrock/app/sssp/sssp_app.cu b/gunrock/app/sssp/sssp_app.cu
diff --git a/gunrock/app/sssp/sssp_enactor.cuh b/gunrock/app/sssp/sssp_enactor.cuh
diff --git a/gunrock/app/sssp/sssp_functor.cuh b/gunrock/app/sssp/sssp_functor.cuh
diff --git a/gunrock/app/sssp/sssp_problem.cuh b/gunrock/app/sssp/sssp_problem.cuh
diff --git a/gunrock/app/topk/topk_app.cu b/gunrock/app/topk/topk_app.cu
diff --git a/gunrock/app/topk/topk_enactor.cuh b/gunrock/app/topk/topk_enactor.cuh
diff --git a/gunrock/app/topk/topk_functor.cuh b/gunrock/app/topk/topk_functor.cuh
diff --git a/gunrock/app/topk/topk_problem.cuh b/gunrock/app/topk/topk_problem.cuh
diff --git a/gunrock/app/wtf/wtf_enactor.cuh b/gunrock/app/wtf/wtf_enactor.cuh
diff --git a/gunrock/app/wtf/wtf_functor.cuh b/gunrock/app/wtf/wtf_functor.cuh
diff --git a/gunrock/app/wtf/wtf_problem.cuh b/gunrock/app/wtf/wtf_problem.cuh
diff --git a/gunrock/coo.cuh b/gunrock/coo.cuh
diff --git a/gunrock/csr.cuh b/gunrock/csr.cuh
diff --git a/gunrock/graphio/market.cuh b/gunrock/graphio/market.cuh
diff --git a/gunrock/graphio/utils.cuh b/gunrock/graphio/utils.cuh
diff --git a/gunrock/gunrock.h b/gunrock/gunrock.h
diff --git a/gunrock/oprtr/advance/kernel.cuh b/gunrock/oprtr/advance/kernel.cuh
@@ -37,6 +37,7 @@ namespace advance {
  * @tparam KernelPolicy Kernel policy type for advance operator.
  * @tparam ProblemData Problem data type for advance operator.
  * @tparam Functor Functor type for the specific problem type.
+ * @tparam Op Operation for gather reduce. mgpu::plus<int> by default.
  *
  * @param[in] d_done                    Pointer of volatile int to the flag to set when we detect incoming frontier is empty
  * @param[in] enactor_stats             EnactorStats object to store enactor related variables and stast
@@ -60,7 +61,14 @@ namespace advance {
  * @param[in] context                   CudaContext pointer for moderngpu APIs
  * @param[in] ADVANCE_TYPE              enumerator of advance type: V2V, V2E, E2V, or E2E
  * @param[in] inverse_graph             whether this iteration of advance operation is in the opposite direction to the previous iteration (false by default)
+ * @param[in] REDUCE_OP                 enumerator of available reduce operations: plus, multiplies, bit_or, bit_and, bit_xor, maximum, minimum. none by default.
+ * @param[in] d_value_to_reduce         array to store values to reduce
+ * @param[out] d_reduce_frontier        neighbor list values for nodes in the output frontier
+ * @param[out] d_reduced_value          array to store reduced values
  */
+
+//TODO: Reduce by neighbor list now only supports LB advance mode.
+
 template <typename KernelPolicy, typename ProblemData, typename Functor>
     void LaunchKernel(
             volatile int                            *d_done,
@@ -84,7 +92,13 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
             util::CtaWorkProgress                   work_progress,
             CudaContext                             &context,
             TYPE                                    ADVANCE_TYPE,
-            bool                                    inverse_graph = false)
+            bool                                    inverse_graph = false,
+            OP                                      REDUCE_OP = KernelPolicy::OP::NONE,
+            REDUCE_TYPE                             R_TYPE = KernelPolicy::REDUCE_TYPE::NONE,
+            typename KernelPolicy::Value            *d_value_to_reduce = NULL,
+            typename KernelPolicy::Value            *d_reduce_frontier = NULL,
+            typename KernelPolicy::Value            *d_reduced_value = NULL)
+
 {
     switch (KernelPolicy::ADVANCE_MODE)
     {
@@ -286,7 +300,10 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
                         work_progress,
                         enactor_stats.advance_kernel_stats,
                         ADVANCE_TYPE,
-                        inverse_graph);
+                        inverse_graph,
+                        R_TYPE,
+                        d_value_to_reduce,
+                        d_reduce_frontier);
             }
             else
             {
@@ -332,10 +349,19 @@ template <typename KernelPolicy, typename ProblemData, typename Functor>
                                         work_progress,
                                         enactor_stats.advance_kernel_stats,
                                         ADVANCE_TYPE,
-                                        inverse_graph);
+                                        inverse_graph,
+                                        R_TYPE,
+                                        d_value_to_reduce,
+                                        d_reduce_frontier);
 
                 //util::DisplayDeviceResults(d_out_key_queue, output_queue_len);
             }
+
+            // TODO:
+            // Transfer d_scanned_edges into exclusive scan result
+            // Do segreduction using d_scanned_edges and d_reduce_frontier
+            // save the result to d_reduced_value
+            //
             break;
         }
         /*case LB:

diff --git a/gunrock/oprtr/advance/kernel_policy.cuh b/gunrock/oprtr/advance/kernel_policy.cuh
@@ -58,6 +58,28 @@ enum TYPE {
     E2E
 };
 
+/**
+ * @brief opeartion to use for mgpu primitives
+ */
+enum OP {
+    NONE,
+    PLUS,
+    MINUS,
+    MULTIPLIES,
+    MODULUS,
+    BIT_OR,
+    BIT_AND,
+    BIT_XOR,
+    MAXIMUM,
+    MINIMUM
+};
+
+enum REDUCE_TYPE {
+    EMPTY,
+    VERTEX,
+    EDGE
+};
+
 /**
  * @brief Kernel configuration policy for all three advance kernels (forward, backward, and load balanced).
  *

diff --git a/gunrock/oprtr/edge_map_backward/cta.cuh b/gunrock/oprtr/edge_map_backward/cta.cuh
diff --git a/gunrock/oprtr/edge_map_backward/kernel.cuh b/gunrock/oprtr/edge_map_backward/kernel.cuh
diff --git a/gunrock/oprtr/edge_map_backward/kernel_policy.cuh b/gunrock/oprtr/edge_map_backward/kernel_policy.cuh
diff --git a/gunrock/oprtr/edge_map_forward/cta.cuh b/gunrock/oprtr/edge_map_forward/cta.cuh
diff --git a/gunrock/oprtr/edge_map_forward/kernel.cuh b/gunrock/oprtr/edge_map_forward/kernel.cuh
diff --git a/gunrock/oprtr/edge_map_forward/kernel_policy.cuh b/gunrock/oprtr/edge_map_forward/kernel_policy.cuh
diff --git a/gunrock/oprtr/edge_map_partitioned/cta.cuh b/gunrock/oprtr/edge_map_partitioned/cta.cuh
diff --git a/gunrock/oprtr/edge_map_partitioned/kernel.cuh b/gunrock/oprtr/edge_map_partitioned/kernel.cuh
@@ -46,6 +46,7 @@ struct Dispatch
 {
     typedef typename KernelPolicy::VertexId VertexId;
     typedef typename KernelPolicy::SizeT    SizeT;
+    typedef typename KernelPolicy::Value    Value;
     typedef typename ProblemData::DataSlice DataSlice;
 
     static __device__ __forceinline__ SizeT GetNeighborListLength(
@@ -100,7 +101,10 @@ struct Dispatch
                                 util::CtaWorkProgress &work_progress,
                                 util::KernelRuntimeStats &kernel_stats,
                                 gunrock::oprtr::advance::TYPE ADVANCE_TYPE,
-                                bool &inverse_graph)
+                                bool &inverse_graph, 
+                                gunrock::oprtr::advance::REDUCE_TYPE R_TYPE,
+                                Value *&d_value_to_reduce,
+                                Value *&d_reduce_frontier)
     {
     }
 
@@ -148,7 +152,10 @@ struct Dispatch
                                 util::CtaWorkProgress &work_progress,
                                 util::KernelRuntimeStats &kernel_stats,
                                 gunrock::oprtr::advance::TYPE ADVANCE_TYPE,
-                                bool &inverse_graph)
+                                bool &inverse_graph, 
+                                gunrock::oprtr::advance::REDUCE_TYPE R_TYPE,
+                                Value *&d_value_to_reduce,
+                                Value *&d_reduce_frontier)
     {
     }
 
@@ -158,6 +165,7 @@ struct Dispatch<KernelPolicy, ProblemData, Functor, true>
 {
     typedef typename KernelPolicy::VertexId         VertexId;
     typedef typename KernelPolicy::SizeT            SizeT;
+    typedef typename KernelPolicy::Value            Value;
     typedef typename ProblemData::DataSlice         DataSlice;
 
     static __device__ __forceinline__ SizeT GetNeighborListLength(
@@ -233,7 +241,11 @@ struct Dispatch<KernelPolicy, ProblemData, Functor, true>
                                 util::CtaWorkProgress &work_progress,
                                 util::KernelRuntimeStats &kernel_stats,
                                 gunrock::oprtr::advance::TYPE &ADVANCE_TYPE,
-                                bool &inverse_graph)
+                                bool &inverse_graph, 
+                                gunrock::oprtr::advance::REDUCE_TYPE R_TYPE,
+                                Value *&d_value_to_reduce,
+                                Value *&d_reduce_frontier)
+
                                 {
                                     if (KernelPolicy::INSTRUMENT && (threadIdx.x == 0 && blockIdx.x == 0)) {
                                         kernel_stats.MarkStart();
@@ -363,6 +375,20 @@ struct Dispatch<KernelPolicy, ProblemData, Functor, true>
                                             SizeT out_index = out_offset+edges_processed+(i-e_offset);
 
                                             {
+                                                if (d_value_to_reduce != NULL) {
+                                                    if (R_TYPE == gunrock::oprtr::advance::VERTEX) {
+                                                    util::io::ModifiedStore<ProblemData::QUEUE_WRITE_MODIFIER>::St(
+                                                                    d_value_to_reduce[u],
+                                                                    d_reduce_frontier + out_index);
+                                                                } else if (R_TYPE == gunrock::oprtr::advance::EDGE) {
+                                                    util::io::ModifiedStore<ProblemData::QUEUE_WRITE_MODIFIER>::St(
+                                                                    d_value_to_reduce[lookup],
+                                                                    d_reduce_frontier + out_index);
+                                                                }
+                                                }
+                                                else if (R_TYPE != gunrock::oprtr::advance::EMPTY) { 
+                                                    // use user-specified function to generate value to reduce
+                                                }
                                                 if (!ProblemData::MARK_PREDECESSORS) {
                                                     if (Functor::CondEdge(label, u, problem, lookup, e_id)) {
                                                         Functor::ApplyEdge(label, u, problem, lookup, e_id);
@@ -565,7 +591,10 @@ struct Dispatch<KernelPolicy, ProblemData, Functor, true>
                                 util::CtaWorkProgress &work_progress,
                                 util::KernelRuntimeStats &kernel_stats,
                                 gunrock::oprtr::advance::TYPE &ADVANCE_TYPE,
-                                bool inverse_graph)
+                                bool inverse_graph, 
+                                gunrock::oprtr::advance::REDUCE_TYPE R_TYPE,
+                                Value *&d_value_to_reduce,
+                                Value *&d_reduce_frontier)
     {
         if (KernelPolicy::INSTRUMENT && (blockIdx.x == 0 && threadIdx.x == 0)) {
             kernel_stats.MarkStart();
@@ -843,7 +872,10 @@ void RelaxPartitionedEdges2(
         util::CtaWorkProgress                   work_progress,
         util::KernelRuntimeStats                kernel_stats,
         gunrock::oprtr::advance::TYPE ADVANCE_TYPE = gunrock::oprtr::advance::V2V,
-        bool                                    inverse_graph = false)
+        bool                                    inverse_graph = false,
+        gunrock::oprtr::advance::REDUCE_TYPE R_TYPE = gunrock::oprtr::advance::EMPTY,
+        typename KernelPolicy::Value            d_value_to_reduce = NULL,
+        typename KernelPolicy::Value            d_reduce_frontier = NULL)
 {
     Dispatch<KernelPolicy, ProblemData, Functor>::RelaxPartitionedEdges2(
             queue_reset,
@@ -867,7 +899,10 @@ void RelaxPartitionedEdges2(
             work_progress,
             kernel_stats,
             ADVANCE_TYPE,
-            inverse_graph);
+            inverse_graph,
+            R_TYPE,
+            d_value_to_reduce,
+            d_reduce_frontier);
 }
 
 /**
@@ -916,7 +951,10 @@ void RelaxLightEdges(
         util::CtaWorkProgress           work_progress,
         util::KernelRuntimeStats        kernel_stats,
         gunrock::oprtr::advance::TYPE ADVANCE_TYPE = gunrock::oprtr::advance::V2V,
-        bool                            inverse_graph = false)
+        bool                            inverse_graph = false, 
+        gunrock::oprtr::advance::REDUCE_TYPE R_TYPE = gunrock::oprtr::advance::EMPTY,
+        typename KernelPolicy::Value    *d_value_to_reduce = NULL,
+        typename KernelPolicy::Value    *d_reduce_frontier = NULL)
 {
     Dispatch<KernelPolicy, ProblemData, Functor>::RelaxLightEdges(
                                 queue_reset,
@@ -937,7 +975,10 @@ void RelaxLightEdges(
                                 work_progress,
                                 kernel_stats,
                                 ADVANCE_TYPE,
-                                inverse_graph);
+                                inverse_graph,
+                                R_TYPE,
+                                d_value_to_reduce,
+                                d_reduce_frontier);
 }
 
 /**

diff --git a/gunrock/oprtr/edge_map_partitioned/kernel_policy.cuh b/gunrock/oprtr/edge_map_partitioned/kernel_policy.cuh
diff --git a/gunrock/oprtr/edge_map_partitioned_backward/kernel.cuh b/gunrock/oprtr/edge_map_partitioned_backward/kernel.cuh
diff --git a/gunrock/oprtr/edge_map_partitioned_backward/kernel_policy.cuh b/gunrock/oprtr/edge_map_partitioned_backward/kernel_policy.cuh
diff --git a/gunrock/oprtr/filter/cta.cuh b/gunrock/oprtr/filter/cta.cuh
diff --git a/gunrock/oprtr/filter/kernel.cuh b/gunrock/oprtr/filter/kernel.cuh
diff --git a/gunrock/oprtr/filter/kernel_policy.cuh b/gunrock/oprtr/filter/kernel_policy.cuh
diff --git a/gunrock/priority_queue/kernel.cuh b/gunrock/priority_queue/kernel.cuh
diff --git a/gunrock/priority_queue/kernel_policy.cuh b/gunrock/priority_queue/kernel_policy.cuh
diff --git a/gunrock/priority_queue/near_far_pile.cuh b/gunrock/priority_queue/near_far_pile.cuh
diff --git a/gunrock/util/basic_utils.h b/gunrock/util/basic_utils.h
diff --git a/gunrock/util/cta_work_distribution.cuh b/gunrock/util/cta_work_distribution.cuh
diff --git a/gunrock/util/cta_work_progress.cuh b/gunrock/util/cta_work_progress.cuh
diff --git a/gunrock/util/cuda_properties.cuh b/gunrock/util/cuda_properties.cuh
diff --git a/gunrock/util/device_intrinsics.cuh b/gunrock/util/device_intrinsics.cuh
diff --git a/gunrock/util/error_utils.cu b/gunrock/util/error_utils.cu
diff --git a/gunrock/util/error_utils.cuh b/gunrock/util/error_utils.cuh
diff --git a/gunrock/util/global_barrier.cuh b/gunrock/util/global_barrier.cuh
diff --git a/gunrock/util/io/initialize_tile.cuh b/gunrock/util/io/initialize_tile.cuh
diff --git a/gunrock/util/io/load_tile.cuh b/gunrock/util/io/load_tile.cuh
diff --git a/gunrock/util/io/modified_load.cuh b/gunrock/util/io/modified_load.cuh
diff --git a/gunrock/util/io/modified_store.cuh b/gunrock/util/io/modified_store.cuh
diff --git a/gunrock/util/io/scatter_tile.cuh b/gunrock/util/io/scatter_tile.cuh
diff --git a/gunrock/util/io/store_tile.cuh b/gunrock/util/io/store_tile.cuh
diff --git a/gunrock/util/kernel_runtime_stats.cuh b/gunrock/util/kernel_runtime_stats.cuh
diff --git a/gunrock/util/mark_segment.cuh b/gunrock/util/mark_segment.cuh
diff --git a/gunrock/util/memset_kernel.cuh b/gunrock/util/memset_kernel.cuh
diff --git a/gunrock/util/multiple_buffering.cuh b/gunrock/util/multiple_buffering.cuh
diff --git a/gunrock/util/numeric_traits.cuh b/gunrock/util/numeric_traits.cuh
diff --git a/gunrock/util/operators.cuh b/gunrock/util/operators.cuh
diff --git a/gunrock/util/random_bits.h b/gunrock/util/random_bits.h
diff --git a/gunrock/util/reduction/cooperative_reduction.cuh b/gunrock/util/reduction/cooperative_reduction.cuh
diff --git a/gunrock/util/reduction/serial_reduce.cuh b/gunrock/util/reduction/serial_reduce.cuh
diff --git a/gunrock/util/reduction/soa/cooperative_soa_reduction.cuh b/gunrock/util/reduction/soa/cooperative_soa_reduction.cuh
diff --git a/gunrock/util/reduction/soa/serial_soa_reduce.cuh b/gunrock/util/reduction/soa/serial_soa_reduce.cuh
diff --git a/gunrock/util/reduction/soa/warp_soa_reduce.cuh b/gunrock/util/reduction/soa/warp_soa_reduce.cuh
diff --git a/gunrock/util/reduction/tree_reduce.cuh b/gunrock/util/reduction/tree_reduce.cuh
diff --git a/gunrock/util/reduction/warp_reduce.cuh b/gunrock/util/reduction/warp_reduce.cuh
diff --git a/gunrock/util/scan/cooperative_scan.cuh b/gunrock/util/scan/cooperative_scan.cuh
diff --git a/gunrock/util/scan/serial_scan.cuh b/gunrock/util/scan/serial_scan.cuh
diff --git a/gunrock/util/scan/soa/cooperative_soa_scan.cuh b/gunrock/util/scan/soa/cooperative_soa_scan.cuh
diff --git a/gunrock/util/scan/soa/serial_soa_scan.cuh b/gunrock/util/scan/soa/serial_soa_scan.cuh
diff --git a/gunrock/util/scan/soa/warp_soa_scan.cuh b/gunrock/util/scan/soa/warp_soa_scan.cuh
diff --git a/gunrock/util/scan/warp_scan.cuh b/gunrock/util/scan/warp_scan.cuh
diff --git a/gunrock/util/select_utils.cuh b/gunrock/util/select_utils.cuh
diff --git a/gunrock/util/soa_tuple.cuh b/gunrock/util/soa_tuple.cuh
diff --git a/gunrock/util/sort_utils.cuh b/gunrock/util/sort_utils.cuh
diff --git a/gunrock/util/srts_details.cuh b/gunrock/util/srts_details.cuh
diff --git a/gunrock/util/srts_grid.cuh b/gunrock/util/srts_grid.cuh
diff --git a/gunrock/util/srts_soa_details.cuh b/gunrock/util/srts_soa_details.cuh
diff --git a/gunrock/util/test_utils.cu b/gunrock/util/test_utils.cu
diff --git a/gunrock/util/test_utils.cuh b/gunrock/util/test_utils.cuh
diff --git a/gunrock/util/test_utils.h b/gunrock/util/test_utils.h
diff --git a/gunrock/util/vector_types.cuh b/gunrock/util/vector_types.cuh
diff --git a/shared_lib_tests/CMakeLists.txt b/shared_lib_tests/CMakeLists.txt
diff --git a/shared_lib_tests/test_bc.c b/shared_lib_tests/test_bc.c
diff --git a/shared_lib_tests/test_bfs.c b/shared_lib_tests/test_bfs.c
diff --git a/shared_lib_tests/test_cc.c b/shared_lib_tests/test_cc.c
diff --git a/shared_lib_tests/test_pr.c b/shared_lib_tests/test_pr.c
diff --git a/shared_lib_tests/test_sssp.c b/shared_lib_tests/test_sssp.c
diff --git a/shared_lib_tests/test_topk.c b/shared_lib_tests/test_topk.c
diff --git a/simple_example/CMakeLists.txt b/simple_example/CMakeLists.txt
diff --git a/simple_example/Makefile b/simple_example/Makefile
diff --git a/simple_example/bips98_606.mtx b/simple_example/bips98_606.mtx
diff --git a/simple_example/cpu_graph_lib.cpp b/simple_example/cpu_graph_lib.cpp
diff --git a/simple_example/cpu_graph_lib.hpp b/simple_example/cpu_graph_lib.hpp
diff --git a/simple_example/simple_example.cu b/simple_example/simple_example.cu
diff --git a/tests/bc/Makefile b/tests/bc/Makefile
diff --git a/tests/bc/ppopp-test.sh b/tests/bc/ppopp-test.sh
diff --git a/tests/bc/run.sh b/tests/bc/run.sh
diff --git a/tests/bc/test_bc.cu b/tests/bc/test_bc.cu
diff --git a/tests/bfs/CMakeLists.txt b/tests/bfs/CMakeLists.txt
diff --git a/tests/bfs/Makefile b/tests/bfs/Makefile
diff --git a/tests/bfs/ppopp-test.sh b/tests/bfs/ppopp-test.sh
diff --git a/tests/bfs/run.sh b/tests/bfs/run.sh
diff --git a/tests/bfs/test_bfs.cu b/tests/bfs/test_bfs.cu
@@ -479,6 +479,7 @@ void RunTests(
     } else if (src_str.compare("largestdegree") == 0) {
         int temp;
         src = graph.GetNodeWithHighestDegree(temp);
+        printf("highest degree:%d\n", temp);
     } else {
         args.GetCmdLineArgument("src", src);
     }
@@ -658,14 +659,17 @@ int main( int argc, char** argv)
 		csr.PrintHistogram();
 
 		// Run tests
-		RunTests(csr, args, *context);
-
+		//RunTests(csr, args, *context);
+    int temp;
+    int src = csr.GetNodeWithHighestDegree(temp);
+    printf("highest degree:%d\n", temp);
     } else {
 
 		// Unknown graph type
 		fprintf(stderr, "Unspecified graph type\n");
 		return 1;
 
 	}
+
 	return 0;
 }
diff --git a/tests/cc/CMakeLists.txt b/tests/cc/CMakeLists.txt
diff --git a/tests/cc/Makefile b/tests/cc/Makefile
diff --git a/tests/cc/ppopp-test.sh b/tests/cc/ppopp-test.sh
diff --git a/tests/cc/run.sh b/tests/cc/run.sh
diff --git a/tests/cc/test_cc.cu b/tests/cc/test_cc.cu
diff --git a/tests/dobfs/Makefile b/tests/dobfs/Makefile
diff --git a/tests/dobfs/ppopp-test.sh b/tests/dobfs/ppopp-test.sh
diff --git a/tests/dobfs/run.sh b/tests/dobfs/run.sh
diff --git a/tests/dobfs/test_dobfs.cu b/tests/dobfs/test_dobfs.cu
diff --git a/tests/hits/Makefile b/tests/hits/Makefile
diff --git a/tests/hits/test_hits.cu b/tests/hits/test_hits.cu
diff --git a/tests/mst/Makefile b/tests/mst/Makefile
diff --git a/tests/mst/test_mst.cu b/tests/mst/test_mst.cu
diff --git a/tests/pr/CMakeLists.txt b/tests/pr/CMakeLists.txt
diff --git a/tests/pr/Makefile b/tests/pr/Makefile
diff --git a/tests/pr/ppopp-test.sh b/tests/pr/ppopp-test.sh
diff --git a/tests/pr/run.sh b/tests/pr/run.sh
diff --git a/tests/pr/test_pr.cu b/tests/pr/test_pr.cu
diff --git a/tests/salsa/Makefile b/tests/salsa/Makefile
diff --git a/tests/salsa/test_salsa.cu b/tests/salsa/test_salsa.cu
diff --git a/tests/sssp/CMakeLists.txt b/tests/sssp/CMakeLists.txt
diff --git a/tests/sssp/Makefile b/tests/sssp/Makefile
diff --git a/tests/sssp/ppopp-test.sh b/tests/sssp/ppopp-test.sh
diff --git a/tests/sssp/run.sh b/tests/sssp/run.sh
diff --git a/tests/sssp/test_sssp.cu b/tests/sssp/test_sssp.cu
diff --git a/tests/topk/Makefile b/tests/topk/Makefile
diff --git a/tests/topk/Makefile_shared_library b/tests/topk/Makefile_shared_library
diff --git a/tests/topk/pre_post_procs/arcs_sample b/tests/topk/pre_post_procs/arcs_sample
diff --git a/tests/topk/pre_post_procs/gpu_result_sample b/tests/topk/pre_post_procs/gpu_result_sample
diff --git a/tests/topk/pre_post_procs/index_sample b/tests/topk/pre_post_procs/index_sample
diff --git a/tests/topk/pre_post_procs/map_sample b/tests/topk/pre_post_procs/map_sample
diff --git a/tests/topk/pre_post_procs/scripts/MapInfoOnly.py b/tests/topk/pre_post_procs/scripts/MapInfoOnly.py
diff --git a/tests/topk/pre_post_procs/scripts/ReadFile.py b/tests/topk/pre_post_procs/scripts/ReadFile.py
diff --git a/tests/topk/pre_post_procs/scripts/ReadMtxOnly.py b/tests/topk/pre_post_procs/scripts/ReadMtxOnly.py
diff --git a/tests/topk/pre_post_procs/scripts/TopK.py b/tests/topk/pre_post_procs/scripts/TopK.py
diff --git a/tests/topk/pre_post_procs/top_k.json b/tests/topk/pre_post_procs/top_k.json
diff --git a/tests/topk/run.sh b/tests/topk/run.sh
diff --git a/tests/topk/test_topk.cu b/tests/topk/test_topk.cu
diff --git a/tests/wtf/Makefile b/tests/wtf/Makefile
diff --git a/tests/wtf/test_wtf.cu b/tests/wtf/test_wtf.cu
diff --git a/wrapper/app/mad/mad_enactor.cuh b/wrapper/app/mad/mad_enactor.cuh
diff --git a/wrapper/app/mad/mad_functor.cuh b/wrapper/app/mad/mad_functor.cuh
diff --git a/wrapper/app/mad/mad_problem.cuh b/wrapper/app/mad/mad_problem.cuh
diff --git a/wrapper/cuda/kernel.cuh b/wrapper/cuda/kernel.cuh
diff --git a/wrapper/cuda/kernel_policy.cuh b/wrapper/cuda/kernel_policy.cuh
diff --git a/wrapper/test/gunrock.cpp b/wrapper/test/gunrock.cpp
diff --git a/wrapper/test/gunrock.h b/wrapper/test/gunrock.h
diff --git a/wrapper/test/gunrock_mad.cu b/wrapper/test/gunrock_mad.cu
diff --git a/wrapper/test/test.cu b/wrapper/test/test.cu
diff --git a/wrapper/test/test_mad.c b/wrapper/test/test_mad.c
diff --git a/wrapper/util/util.cu b/wrapper/util/util.cu
diff --git a/wrapper/util/util.h b/wrapper/util/util.h