Merge pull request #103 from gunrock/dev

Bring changes over, so I can improve kernels.
gunrock · Dec 24, 2021 · 93b99fc · 93b99fc
2 parents 7939a31 + 559938a
commit 93b99fc
Show file tree

Hide file tree

Showing 14 changed files with 340 additions and 118 deletions.
diff --git a/CMakeLists.txt b/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)
diff --git a/benchmarks/CMakeLists.txt b/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()
diff --git a/benchmarks/for.cu b/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
diff --git a/cmake/FetchCUB.cmake b/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})

diff --git a/cmake/FetchModernGPU.cmake b/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})

diff --git a/cmake/FetchNVBench.cmake b/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})
diff --git a/cmake/FetchThrustCUB.cmake b/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})

diff --git a/include/gunrock/container/vector.hxx b/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

diff --git a/include/gunrock/cuda/detail/launch_box.hxx b/include/gunrock/cuda/detail/launch_box.hxx
@@ -22,13 +22,12 @@ namespace detail {
  * @brief Abstract base class for launch parameters.
  *
  * @tparam sm_flags_ Bitwise flags indicating SM versions (`sm_flag_t` enum).
+ * @tparam shared_memory_bytes_ Number of bytes of shared memory to allocate.
  */
-template <sm_flag_t sm_flags_>
+template <sm_flag_t sm_flags_, size_t shared_memory_bytes_>
 struct launch_params_base_t {
-  enum : unsigned { sm_flags = sm_flags_ };
-
- protected:
-  launch_params_base_t() {}
+  static constexpr sm_flag_t sm_flags = sm_flags_;
+  static constexpr size_t shared_memory_bytes = shared_memory_bytes_;
 };
 
 /**
@@ -73,7 +72,8 @@ struct raise_not_found_error_t {
  * was inspired by this Stack Overflow solution:
  * https://stackoverflow.com/a/67155114/13232647.
  *
- * @tparam lp_v Pack of `launch_params_t` types for each desired arch.
+ * @tparam lp_v Pack of `launch_params_t` types for each corresponding
+ * architecture(s).
  */
 template <typename... lp_v>
 using match_launch_params_t = decltype(std::tuple_cat(