Merge pull request #99 from gunrock/bench

Benchmarking support, cmake version++, and sample csr()

CMakeLists.txt

@@ -1,5 +1,5 @@
 # We can lower this if needed.
-cmake_minimum_required(VERSION 3.19 FATAL_ERROR)
+cmake_minimum_required(VERSION 3.20.1 FATAL_ERROR)
 
 # begin /* Update Essentials version */
 set(ESSENTIALS_VERSION_MAJOR 2)
@@ -77,6 +77,8 @@ target_compile_definitions(essentials
     ESSENTIALS_VERSION=${ESSENTIALS_VERSION}
 )
 
+message(STATUS "Essentials CUDA Architecture: ${ESSENTIALS_ARCHITECTURES}")
+
 ####################################################
 ############ TARGET COMPILE FEATURES ###############
 ####################################################
@@ -175,3 +177,25 @@ option(ESSENTIALS_BUILD_TESTS
 if(ESSENTIALS_BUILD_TESTS)
   add_subdirectory(unittests)
 endif(ESSENTIALS_BUILD_TESTS)
+
+####################################################
+################ BUILD BENCHMARKS  #################
+####################################################
+option(ESSENTIALS_BUILD_BENCHMARKS
+  "If on, builds essentials with benchmarking support."
+  OFF)
+
+
+
+# Subdirectories for examples, testing and documentation
+if(ESSENTIALS_BUILD_BENCHMARKS)
+  # ... see https://github.com/NVIDIA/nvbench/issues/66
+  set(NVBench_ENABLE_NVML OFF)
+  # ... set cuda architecture for nvbench.
+  set(CMAKE_CUDA_ARCHITECTURES ${ESSENTIALS_ARCHITECTURES})
+  include(${PROJECT_SOURCE_DIR}/cmake/FetchNVBench.cmake)
+  target_link_libraries(essentials 
+    INTERFACE nvbench::main
+  )
+  add_subdirectory(benchmarks)
+endif(ESSENTIALS_BUILD_BENCHMARKS)

benchmarks/CMakeLists.txt

@@ -0,0 +1,20 @@
+set(BENCHMARK_SOURCES
+  for.cu
+)
+
+foreach(SOURCE IN LISTS BENCHMARK_SOURCES)
+  get_filename_component(BENCHMARK_NAME ${SOURCE} NAME_WLE)
+  add_executable(${BENCHMARK_NAME} ${SOURCE})
+  target_link_libraries(${BENCHMARK_NAME} 
+    PRIVATE essentials
+    PRIVATE nvbench::main
+  )
+  get_target_property(ESSENTIALS_ARCHITECTURES 
+    essentials CUDA_ARCHITECTURES
+  )
+  set_target_properties(${BENCHMARK_NAME} 
+    PROPERTIES 
+      CUDA_ARCHITECTURES ${ESSENTIALS_ARCHITECTURES}
+  )
+  message(STATUS "Benchmark Added: ${BENCHMARK_NAME}")
+endforeach()

benchmarks/for.cu

@@ -0,0 +1,48 @@
+#include <gunrock/error.hxx>
+#include <gunrock/graph/graph.hxx>
+#include <gunrock/formats/formats.hxx>
+#include <gunrock/cuda/cuda.hxx>
+#include <gunrock/framework/operators/for/for.hxx>
+#include <gunrock/util/sample.hxx>
+
+#include <nvbench/nvbench.cuh>
+
+namespace gunrock {
+namespace benchmark {
+void parallel_for(nvbench::state& state) {
+  auto csr = sample::csr();
+
+  // --
+  // Build graph
+
+  auto G = graph::build::from_csr<memory_space_t::device, graph::view_t::csr>(
+      csr.number_of_rows,               // rows
+      csr.number_of_columns,            // columns
+      csr.number_of_nonzeros,           // nonzeros
+      csr.row_offsets.data().get(),     // row_offsets
+      csr.column_indices.data().get(),  // column_indices
+      csr.nonzero_values.data().get()   // values
+  );
+
+  // Initialize the context.
+  cuda::device_id_t device = 0;
+  cuda::multi_context_t context(device);
+
+  vector_t<int> vertices(G.get_number_of_vertices());
+  auto d_vertices = vertices.data().get();
+
+  auto f = [=] __device__(int const& v) -> void { d_vertices[v] = v; };
+
+  state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
+    operators::parallel_for::execute<operators::parallel_for_each_t::vertex>(
+        G,       // graph
+        f,       // lambda function
+        context  // context
+    );
+  });
+}
+
+NVBENCH_BENCH(parallel_for).set_name("parallel_for");
+
+}  // namespace benchmark
+}  // namespace gunrock

cmake/FetchCUB.cmake

@@ -1,7 +1,7 @@
 include(FetchContent)
 set(FETCHCONTENT_QUIET ON)
 
-message("-- Cloning External Project: CUB")
+message(STATUS "Cloning External Project: CUB")
 get_filename_component(FC_BASE "../externals"
                 REALPATH BASE_DIR "${CMAKE_BINARY_DIR}")
 set(FETCHCONTENT_BASE_DIR ${FC_BASE})

cmake/FetchModernGPU.cmake

@@ -1,7 +1,7 @@
 include(FetchContent)
 set(FETCHCONTENT_QUIET ON)
 
-message("-- Cloning External Project: ModernGPU")
+message(STATUS "Cloning External Project: ModernGPU")
 get_filename_component(FC_BASE "../externals"
                 REALPATH BASE_DIR "${CMAKE_BINARY_DIR}")
 set(FETCHCONTENT_BASE_DIR ${FC_BASE})

cmake/FetchNVBench.cmake

@@ -0,0 +1,27 @@
+include(FetchContent)
+set(FETCHCONTENT_QUIET ON)
+
+message(STATUS "Cloning External Project: NVBench")
+get_filename_component(FC_BASE "../externals"
+                REALPATH BASE_DIR "${CMAKE_BINARY_DIR}")
+set(FETCHCONTENT_BASE_DIR ${FC_BASE})
+
+FetchContent_Declare(
+    nvbench
+    GIT_REPOSITORY https://github.com/NVIDIA/nvbench.git
+    GIT_TAG        main
+)
+
+FetchContent_GetProperties(nvbench)
+if(NOT nvbench_POPULATED)
+  FetchContent_Populate(
+    nvbench
+  )
+endif()
+
+# Exposing nvbench's source and include directory
+set(NVBENCH_INCLUDE_DIR "${nvbench_SOURCE_DIR}")
+set(NVBENCH_BUILD_DIR "${nvbench_BINARY_DIR}")
+
+# Add subdirectory ::nvbench
+add_subdirectory(${NVBENCH_INCLUDE_DIR} ${NVBENCH_BUILD_DIR})

cmake/FetchThrustCUB.cmake

@@ -1,7 +1,7 @@
 include(FetchContent)
 set(FETCHCONTENT_QUIET ON)
 
-message("-- Cloning External Project: Thrust")
+message(STATUS "Cloning External Project: Thrust")
 get_filename_component(FC_BASE "../externals"
                 REALPATH BASE_DIR "${CMAKE_BINARY_DIR}")
 set(FETCHCONTENT_BASE_DIR ${FC_BASE})

include/gunrock/container/vector.hxx

@@ -24,7 +24,7 @@ namespace gunrock {
 
 using namespace memory;
 
-template <typename type_t, memory_space_t space>
+template <typename type_t, memory_space_t space = memory_space_t::device>
 using vector_t =
     std::conditional_t<space == memory_space_t::host,  // condition
                        thrust::host_vector<type_t>,    // host_type

include/gunrock/formats/csr.hxx

@@ -49,6 +49,19 @@ struct csr_t {
 
   ~csr_t() {}
 
+  /**
+   * @brief Copy constructor.
+   * @param rhs
+   */
+  template <typename _csr_t>
+  csr_t(const _csr_t& rhs)
+      : number_of_rows(rhs.number_of_rows),
+        number_of_columns(rhs.number_of_columns),
+        number_of_nonzeros(rhs.number_of_nonzeros),
+        row_offsets(rhs.row_offsets),
+        column_indices(rhs.column_indices),
+        nonzero_values(rhs.nonzero_values) {}
+
   /**
    * @brief Convert a Coordinate Sparse Format into Compressed Sparse Row
    * Format.
@@ -163,12 +176,12 @@ struct csr_t {
       thrust::host_vector<index_t> h_column_indices(number_of_nonzeros);
       thrust::host_vector<value_t> h_nonzero_values(number_of_nonzeros);
 
-      assert(fread(memory::raw_pointer_cast(h_row_offsets.data()), sizeof(offset_t),
-            number_of_rows + 1, file) != 0);
-      assert(fread(memory::raw_pointer_cast(h_column_indices.data()), sizeof(index_t),
-            number_of_nonzeros, file) != 0);
-      assert(fread(memory::raw_pointer_cast(h_nonzero_values.data()), sizeof(value_t),
-            number_of_nonzeros, file) != 0);
+      assert(fread(memory::raw_pointer_cast(h_row_offsets.data()),
+                   sizeof(offset_t), number_of_rows + 1, file) != 0);
+      assert(fread(memory::raw_pointer_cast(h_column_indices.data()),
+                   sizeof(index_t), number_of_nonzeros, file) != 0);
+      assert(fread(memory::raw_pointer_cast(h_nonzero_values.data()),
+                   sizeof(value_t), number_of_nonzeros, file) != 0);
 
       // Copy data from host to device
       row_offsets = h_row_offsets;
@@ -178,12 +191,12 @@ struct csr_t {
     } else {
       assert(space == memory_space_t::host);
 
-      assert(fread(memory::raw_pointer_cast(row_offsets.data()), sizeof(offset_t),
-            number_of_rows + 1, file) != 0);
-      assert(fread(memory::raw_pointer_cast(column_indices.data()), sizeof(index_t),
-            number_of_nonzeros, file) != 0);
-      assert(fread(memory::raw_pointer_cast(nonzero_values.data()), sizeof(value_t),
-            number_of_nonzeros, file) != 0);
+      assert(fread(memory::raw_pointer_cast(row_offsets.data()),
+                   sizeof(offset_t), number_of_rows + 1, file) != 0);
+      assert(fread(memory::raw_pointer_cast(column_indices.data()),
+                   sizeof(index_t), number_of_nonzeros, file) != 0);
+      assert(fread(memory::raw_pointer_cast(nonzero_values.data()),
+                   sizeof(value_t), number_of_nonzeros, file) != 0);
     }
   }
 

include/gunrock/util/sample.hxx

@@ -0,0 +1,76 @@
+/**
+ * @file sample.hxx
+ * @author Muhammad Osama (mosama@ucdavis.edu)
+ * @brief
+ * @version 0.1
+ * @date 2021-12-23
+ *
+ * @copyright Copyright (c) 2021
+ *
+ */
+#pragma once
+
+#include <gunrock/formats/formats.hxx>
+
+namespace gunrock {
+namespace sample {
+
+using namespace memory;
+
+template <memory_space_t space = memory_space_t::device,
+          typename vertex_t = int,
+          typename edge_t = int,
+          typename weight_t = float>
+format::csr_t<space, vertex_t, edge_t, weight_t> csr() {
+  // Logical Matrix Representation
+  // r/c  0 1 2 3
+  //  0 [ 0 0 0 0 ]
+  //  1 [ 5 8 0 0 ]
+  //  2 [ 0 0 3 0 ]
+  //  3 [ 0 6 0 0 ]
+
+  // Logical Graph Representation
+  // (i, j) [w]
+  // (1, 0) [5]
+  // (1, 1) [8]
+  // (2, 2) [3]
+  // (3, 1) [6]
+
+  // CSR Matrix Representation
+  // V            = [ 5 8 3 6 ]
+  // COL_INDEX    = [ 0 1 2 1 ]
+  // ROW_OFFSETS  = [ 0 0 2 3 4 ]
+  using csr_t = format::csr_t<memory_space_t::host, vertex_t, edge_t, weight_t>;
+  csr_t matrix(4, 4, 4);
+
+  // Row Offsets
+  matrix.row_offsets[0] = 0;
+  matrix.row_offsets[1] = 0;
+  matrix.row_offsets[2] = 2;
+  matrix.row_offsets[3] = 3;
+  matrix.row_offsets[4] = 4;
+
+  // Column Indices
+  matrix.column_indices[0] = 0;
+  matrix.column_indices[1] = 1;
+  matrix.column_indices[2] = 2;
+  matrix.column_indices[3] = 3;
+
+  // Non-zero values
+  matrix.nonzero_values[0] = 5;
+  matrix.nonzero_values[1] = 8;
+  matrix.nonzero_values[2] = 3;
+  matrix.nonzero_values[3] = 6;
+
+  if (space == memory_space_t::host) {
+    return matrix;
+  } else {
+    using d_csr_t =
+        format::csr_t<memory_space_t::device, vertex_t, edge_t, weight_t>;
+    d_csr_t d_matrix(matrix);
+    return d_matrix;
+  }
+}
+
+}  // namespace sample
+}  // namespace gunrock