Color CPU reference + correctness checking

examples/color/color.cu

@@ -1,6 +1,7 @@
 #include <set>
 
 #include <gunrock/applications/color.hxx>
+#include "color_cpu.hxx"  // Reference implementation
 
 using namespace gunrock;
 using namespace memory;
@@ -24,7 +25,10 @@ void test_color(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));
 
   // --
@@ -46,17 +50,34 @@ void test_color(int num_arguments, char** argument_array) {
   thrust::device_vector<vertex_t> colors(n_vertices);
 
   // --
-  // Run problem
+  // GPU Run
 
   float gpu_elapsed = gunrock::color::run(G, colors.data().get());
+  std::cout << "GPU Elapsed Time : " << gpu_elapsed << " (ms)" << std::endl;
+
+  // --
+  // CPU Run
+
+  thrust::host_vector<vertex_t> h_colors(n_vertices);
+
+  float cpu_elapsed = color_cpu::run<csr_t, vertex_t, edge_t, weight_t>(
+      csr, h_colors.data());
+
+  int n_errors = color_cpu::compute_error<csr_t, vertex_t, edge_t, weight_t>(
+    csr, colors, h_colors);
 
   // --
   // Log
 
-  std::cout << "colors[:40] = ";
+  std::cout << "GPU colors[:40] = ";
   gunrock::print::head<weight_t>(colors, 40);
-  std::cout << std::endl;
-  std::cout << "GPU Elapsed Time: " << gpu_elapsed << " (ms)" << std::endl;
+
+  std::cout << "CPU colors[:40] = ";
+  gunrock::print::head<weight_t>(h_colors, 40);
+
+  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/color/color_cpu.hxx

@@ -0,0 +1,91 @@
+#pragma once
+
+#include <chrono>
+
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/random.h>
+
+#include <gunrock/algorithms/generate/random.hxx>
+
+namespace color_cpu {
+
+using namespace std;
+using namespace std::chrono;
+
+template <typename csr_t, typename vertex_t, typename edge_t, typename weight_t>
+float run(csr_t& csr, vertex_t* colors) {
+
+  thrust::host_vector<edge_t> row_offsets(csr.row_offsets);         // Copy data to CPU
+  thrust::host_vector<vertex_t> column_indices(csr.column_indices);
+
+  auto t_start = high_resolution_clock::now();
+
+  vertex_t n_vertices = csr.number_of_rows;
+  for (vertex_t i = 0; i < n_vertices; i++) colors[i] = -1;
+
+  thrust::host_vector<vertex_t> randoms(n_vertices);
+  gunrock::algo::generate::random::uniform_distribution(0, n_vertices, randoms.begin());
+
+  int color  = 0;
+  int n_left = n_vertices;
+  while(n_left > 0) {
+    for(vertex_t v = 0; v < n_vertices; v++) {
+      if(colors[v] != -1) continue;
+
+      edge_t start_edge    = row_offsets[v];
+      edge_t num_neighbors = row_offsets[v + 1] - row_offsets[v];
+
+      bool colormax = true;
+      bool colormin = true;
+
+      for (edge_t e = start_edge; e < start_edge + num_neighbors; ++e) {
+        vertex_t u = column_indices[e];
+
+        if ((colors[u] != -1) && (colors[u] != color + 1) && (colors[u] != color + 2) || (v == u))
+          continue;
+
+        if (randoms[v] <= randoms[u]) colormax = false;
+        if (randoms[v] >= randoms[u]) colormin = false;
+      }
+
+      if (colormax) {
+        colors[v] = color + 1;
+      } else if (colormin) {
+        colors[v] = color + 2;
+      }
+
+      if(colormax || colormin) n_left--;
+    }
+
+    color += 2;
+  }
+
+  auto t_stop = high_resolution_clock::now();
+  auto elapsed = duration_cast<microseconds>(t_stop - t_start).count();
+  return (float)elapsed / 1000;
+}
+
+template <typename csr_t, typename vertex_t, typename edge_t, typename weight_t>
+float compute_error(csr_t& csr, thrust::device_vector<vertex_t> _gpu_colors, thrust::host_vector<vertex_t> cpu_colors) {
+
+  thrust::host_vector<edge_t>   row_offsets(csr.row_offsets);       // Copy data to CPU
+  thrust::host_vector<vertex_t> column_indices(csr.column_indices);
+  thrust::host_vector<vertex_t> gpu_colors(_gpu_colors);
+
+  vertex_t n_vertices = csr.number_of_rows;
+
+  vertex_t cpu_errors = 0;
+  vertex_t gpu_errors = 0;
+
+  for(vertex_t v = 0; v < n_vertices; v++) {
+    for(edge_t e = row_offsets[v]; e < row_offsets[v + 1]; e++) {
+      vertex_t u = column_indices[e];
+      if(gpu_colors[u] == gpu_colors[v] || gpu_colors[v] == -1) gpu_errors++;
+      // if(cpu_colors[u] == cpu_colors[v] || cpu_colors[v] == -1) cpu_errors++;
+    }
+  }
+
+  return gpu_errors + cpu_errors;
+}
+
+}  // namespace sssp_cpu

include/gunrock/algorithms/generate/random.hxx

@@ -13,7 +13,7 @@ template <typename index_t, typename iterator_t>
 void uniform_distribution(index_t begin, index_t end, iterator_t input) {
   using type_t = typename std::iterator_traits<iterator_t>::value_type;
 
-  auto generate_random = [] __device__(int i) -> type_t {
+  auto generate_random = [] __host__ __device__(int i) -> type_t {
     thrust::default_random_engine rng;
     /* thrust::uniform_real_distribution<type_t> uni_dist; */
     rng.discard(i);