Merge pull request #56 from gunrock/master

Merging changes to `dev` branch.

examples/CMakeLists.txt

@@ -4,4 +4,5 @@ add_subdirectory(bfs)
 add_subdirectory(color)
 add_subdirectory(geo)
 add_subdirectory(pr)
+add_subdirectory(mtx2bin)
 # end /* Add examples' subdirectories */

examples/bfs/bfs.cu

@@ -1,4 +1,5 @@
 #include <gunrock/applications/bfs.hxx>
+#include "bfs_cpu.hxx"  // Reference implementation
 
 using namespace gunrock;
 using namespace memory;
@@ -22,8 +23,12 @@ void test_bfs(int num_arguments, char** argument_array) {
   std::string filename = argument_array[1];
 
   io::matrix_market_t<vertex_t, edge_t, weight_t> mm;
-  format::csr_t<memory_space_t::device, vertex_t, edge_t, weight_t> csr;
+
+  using csr_t =
+      format::csr_t<memory_space_t::device, vertex_t, edge_t, weight_t>;
+  csr_t csr;
   csr.from_coo(mm.load(filename));
+
   thrust::device_vector<vertex_t> row_indices(csr.number_of_nonzeros);
   thrust::device_vector<edge_t> column_offsets(csr.number_of_columns + 1);
 
@@ -54,17 +59,42 @@ void test_bfs(int num_arguments, char** argument_array) {
   // --
   // Run problem
 
-  float elapsed = gunrock::bfs::run(G, single_source, distances.data().get(),
+  float gpu_elapsed = gunrock::bfs::run(G, single_source, distances.data().get(),
                                     predecessors.data().get());
 
+  // --
+  // CPU Run
+
+  thrust::host_vector<vertex_t> h_distances(n_vertices);
+  thrust::host_vector<vertex_t> h_predecessors(n_vertices);
+
+  float cpu_elapsed = bfs_cpu::run<csr_t, vertex_t, edge_t>(
+      csr, single_source, h_distances.data(), h_predecessors.data());
+
+  int n_errors = bfs_cpu::compute_error(distances, h_distances);
+
   // --
   // Log
 
-  std::cout << "Distances (output) = ";
-  thrust::copy(distances.begin(), distances.end(),
+  std::cout << "GPU Distances (output) = ";
+  thrust::copy(distances.begin(),
+               (distances.size() < 40) ? distances.begin() + distances.size()
+                                       : distances.begin() + 40,
                std::ostream_iterator<vertex_t>(std::cout, " "));
   std::cout << std::endl;
-  std::cout << "BFS Elapsed Time: " << elapsed << " (ms)" << std::endl;
+
+  std::cout << "CPU Distances (output) = ";
+  thrust::copy(h_distances.begin(),
+               (h_distances.size() < 40)
+                   ? h_distances.begin() + h_distances.size()
+                   : h_distances.begin() + 40,
+               std::ostream_iterator<vertex_t>(std::cout, " "));
+  std::cout << std::endl;
+
+  std::cout << "GPU Elapsed Time : " << gpu_elapsed << " (ms)" << std::endl;
+  std::cout << "CPU Elapsed Time : " << cpu_elapsed << " (ms)" << std::endl;
+  std::cout << "Number of errors : " << n_errors << std::endl;
+
 }
 
 int main(int argc, char** argv) {

examples/bfs/bfs_cpu.hxx

@@ -0,0 +1,80 @@
+#pragma once
+
+#include <chrono>
+#include <vector>
+#include <queue>
+
+namespace bfs_cpu {
+
+using namespace std;
+using namespace std::chrono;
+
+template <typename vertex_t>
+class prioritize {
+ public:
+  bool operator()(pair<vertex_t, vertex_t>& p1, pair<vertex_t, vertex_t>& p2) {
+    return p1.second > p2.second;
+  }
+};
+
+template <typename csr_t, typename vertex_t, typename edge_t>
+float run(csr_t& csr,
+          vertex_t& single_source,
+          vertex_t* distances,
+          vertex_t* predecessors) {
+  thrust::host_vector<edge_t> row_offsets(csr.row_offsets);  // Copy data to CPU
+  thrust::host_vector<vertex_t> column_indices(csr.column_indices);
+
+  for (vertex_t i = 0; i < csr.number_of_rows; i++)
+    distances[i] = std::numeric_limits<vertex_t>::max();
+
+  auto t_start = high_resolution_clock::now();
+
+  distances[single_source] = 0;
+
+  priority_queue<pair<vertex_t, vertex_t>,
+                 std::vector<pair<vertex_t, vertex_t>>,
+                 prioritize<vertex_t>> pq;
+
+  pq.push(make_pair(single_source, 0.0));
+
+  while (!pq.empty()) {
+    pair<vertex_t, vertex_t> curr = pq.top();
+    pq.pop();
+
+    vertex_t curr_node = curr.first;
+    vertex_t curr_dist = curr.second;
+
+    vertex_t start = row_offsets[curr_node];
+    vertex_t end = row_offsets[curr_node + 1];
+
+    for (vertex_t offset = start; offset < end; offset++) {
+      vertex_t neib = column_indices[offset];
+      vertex_t new_dist = curr_dist + 1;
+      if (new_dist < distances[neib]) {
+        distances[neib] = new_dist;
+        pq.push(make_pair(neib, new_dist));
+      }
+    }
+  }
+
+  auto t_stop = high_resolution_clock::now();
+  auto elapsed = duration_cast<microseconds>(t_stop - t_start).count();
+  return (float)elapsed / 1000;
+}
+
+template <typename val_t>
+int compute_error(thrust::device_vector<val_t> _gpu_result,
+                  thrust::host_vector<val_t> cpu_result) {
+  thrust::host_vector<val_t> gpu_result(_gpu_result);
+
+  int n_errors = 0;
+  for (int i = 0; i < cpu_result.size(); i++) {
+    if (gpu_result[i] != cpu_result[i]) {
+      n_errors++;
+    }
+  }
+  return n_errors;
+}
+
+}  // namespace bfs_cpu

examples/color/color.cu

@@ -48,16 +48,15 @@ void test_color(int num_arguments, char** argument_array) {
   // --
   // Run problem
 
-  float elapsed = gunrock::color::run(G, colors.data().get());
+  float gpu_elapsed = gunrock::color::run(G, colors.data().get());
 
   // --
   // Log
 
-  std::cout << "Colors (output) = ";
-  thrust::copy(colors.begin(), colors.end(),
-               std::ostream_iterator<weight_t>(std::cout, " "));
+  std::cout << "colors[:40] = ";
+  gunrock::print::head<weight_t>(colors, 40);
   std::cout << std::endl;
-  std::cout << "color Elapsed Time: " << elapsed << " (ms)" << std::endl;
+  std::cout << "GPU Elapsed Time: " << gpu_elapsed << " (ms)" << std::endl;
 }
 
 int main(int argc, char** argv) {

examples/mtx2bin/CMakeLists.txt

@@ -0,0 +1,21 @@
+# begin /* Set the application name. */
+set(APPLICATION_NAME mtx2bin)
+# end /* Set the application name. */
+
+# begin /* Add CUDA executables */
+add_executable(${APPLICATION_NAME})
+
+set(SOURCE_LIST 
+    ${APPLICATION_NAME}.cu
+)
+
+target_sources(${APPLICATION_NAME} PRIVATE ${SOURCE_LIST})
+target_link_libraries(${APPLICATION_NAME} PRIVATE essentials)
+get_target_property(ESSENTIALS_ARCHITECTURES essentials CUDA_ARCHITECTURES)
+set_target_properties(${APPLICATION_NAME} 
+    PROPERTIES 
+        CUDA_ARCHITECTURES ${ESSENTIALS_ARCHITECTURES}
+) # XXX: Find a better way to inherit essentials properties.
+
+message("-- Example Added: ${APPLICATION_NAME}")
+# end /* Add CUDA executables */

examples/mtx2bin/mtx2bin.cu

@@ -0,0 +1,44 @@
+#include <cstdlib>  // EXIT_SUCCESS
+
+#include <gunrock/applications/application.hxx>
+
+using namespace gunrock;
+using namespace memory;
+
+void test_mtx2bin(int num_arguments, char** argument_array) {
+  if (num_arguments != 2) {
+    std::cerr << "usage: ./bin/mtx2bin <inpath>" << std::endl;
+    exit(1);
+  }
+
+  // --
+  // Define types
+
+  using vertex_t = int;
+  using edge_t = int;
+  using weight_t = float;
+
+  // --
+  // IO
+
+  std::string inpath  = argument_array[1];
+  std::string outpath = inpath + ".csr";
+
+  io::matrix_market_t<vertex_t, edge_t, weight_t> mm;
+
+  using csr_t = format::csr_t<memory::memory_space_t::device, vertex_t, edge_t, weight_t>;
+  csr_t csr;
+  csr.from_coo(mm.load(inpath));
+
+  std::cout << "csr.number_of_rows     = " << csr.number_of_rows     << std::endl;
+  std::cout << "csr.number_of_columns  = " << csr.number_of_columns  << std::endl;
+  std::cout << "csr.number_of_nonzeros = " << csr.number_of_nonzeros << std::endl;
+  std::cout << "writing to             = " << outpath                << std::endl;
+
+  csr.write_binary(outpath);
+}
+
+int main(int argc, char** argv) {
+  test_mtx2bin(argc, argv);
+  return EXIT_SUCCESS;
+}

examples/pr/pr.cu

@@ -16,18 +16,24 @@ void test_sssp(int num_arguments, char** argument_array) {
   using edge_t = int;
   using weight_t = float;
 
+  using csr_t = format::csr_t<memory_space_t::device, vertex_t, edge_t, weight_t>;
+  csr_t csr;
+
   // --
   // IO
 
   std::string filename = argument_array[1];
-
-  io::matrix_market_t<vertex_t, edge_t, weight_t> mm;
-
-  using csr_t =
-      format::csr_t<memory_space_t::device, vertex_t, edge_t, weight_t>;
-  csr_t csr;
-  csr.from_coo(mm.load(filename));
-
+
+  if(util::is_market(filename)) {
+    io::matrix_market_t<vertex_t, edge_t, weight_t> mm;
+    csr.from_coo(mm.load(filename));
+  } else if(util::is_binary_csr(filename)) {
+    csr.read_binary(filename);
+  } else {
+    std::cerr << "Unknown file format: " << filename << std::endl;
+    exit(1);
+  }
+
   // --
   // Build graph
 
@@ -54,19 +60,13 @@ void test_sssp(int num_arguments, char** argument_array) {
   // --
   // GPU Run
 
-  std::cout << "gunrock::pr::run -- starting" << std::endl;
   float gpu_elapsed = gunrock::pr::run(G, alpha, tol, p.data().get());
-  std::cout << "gunrock::pr::run -- complete" << std::endl;
 
   // --
   // Log + Validate
 
-  std::cout << "GPU p (output) = ";
-  thrust::copy(p.begin(),
-               (p.size() < 40) ? p.begin() + p.size()
-                                       : p.begin() + 40,
-               std::ostream_iterator<weight_t>(std::cout, " "));
-  std::cout << std::endl;
+  std::cout << "GPU p[:40] = ";
+  gunrock::print::head<weight_t>(p, 40);
   std::cout << "GPU Elapsed Time : " << gpu_elapsed << " (ms)" << std::endl;
 }
 

examples/sssp/sssp.cu

@@ -71,20 +71,11 @@ void test_sssp(int num_arguments, char** argument_array) {
   // --
   // Log + Validate
 
-  std::cout << "GPU Distances (output) = ";
-  thrust::copy(distances.begin(),
-               (distances.size() < 40) ? distances.begin() + distances.size()
-                                       : distances.begin() + 40,
-               std::ostream_iterator<weight_t>(std::cout, " "));
-  std::cout << std::endl;
+  std::cout << "GPU distances[:40] = ";
+  gunrock::print::head<weight_t>(distances, 40);
 
   std::cout << "CPU Distances (output) = ";
-  thrust::copy(h_distances.begin(),
-               (h_distances.size() < 40)
-                   ? h_distances.begin() + h_distances.size()
-                   : h_distances.begin() + 40,
-               std::ostream_iterator<weight_t>(std::cout, " "));
-  std::cout << std::endl;
+  gunrock::print::head<weight_t>(h_distances, 40);
 
   std::cout << "GPU Elapsed Time : " << gpu_elapsed << " (ms)" << std::endl;
   std::cout << "CPU Elapsed Time : " << cpu_elapsed << " (ms)" << std::endl;

include/gunrock/applications/application.hxx

@@ -19,6 +19,7 @@
 
 // Utility includes
 #include <gunrock/util/math.hxx>
+#include <gunrock/util/print.hxx>
 
 // Format includes
 #include <gunrock/formats/formats.hxx>