Merge tune and benchmark variables non intrusively

This piece of code integrates and shows how to tune some variables,
namely the multiple used to decide on the number of warps used in COO.
This is a suggestion of what I could find as the simplest and most versatile
solution to this issue. If anyone has another idea I would be happy to try it out.

How: This is integrated into the SPMV benchmarks for now and is entirely
controlled with the new CMake option `GINKGO_BENCHMARK_ENABLE_TUNING`.

The setup works through the use of (dynamic) global variables so that
tuning is as little intrusive as possible. This should work even with the
"implementation selection" since this also has a dynamic component which
could use the same strategy as shown here.

Caveat: With this setup, only benchmarks can be compiled, other
executable codes such as tests would require to also add
`benchmark/utils/tuning_variables.cpp` or equivalent. This option should
of course not be turned on for general Ginkgo usage, but only for
tuning.

Related PR: #692

CMakeLists.txt

@@ -44,6 +44,9 @@ option(GINKGO_CHECK_CIRCULAR_DEPS
     OFF)
 option(GINKGO_CONFIG_LOG_DETAILED
     "Enable printing of detailed configuration log to screen in addition to the writing of files," OFF)
+option(GINKGO_BENCHMARK_ENABLE_TUNING
+    "Enable tuning variables in the benchmarks. For specific use cases, manual code changes could be required."
+    OFF)
 set(GINKGO_VERBOSE_LEVEL "1" CACHE STRING
     "Verbosity level. Put 0 to turn off. 1 activates a few important messages.")
 if(MSVC)
@@ -74,6 +77,13 @@ option(GINKGO_BUILD_HWLOC "Build Ginkgo with HWLOC. Default is ON. If a system H
 
 set(GINKGO_CIRCULAR_DEPS_FLAGS "-Wl,--no-undefined")
 
+if(GINKGO_BENCHMARK_ENABLE_TUNING)
+    # In this state, the tests and examples cannot be compiled without extra
+    # complexity/intrusiveness, so we simply disable them.
+    set(GINKGO_BUILD_TESTS OFF)
+    set(GINKGO_BUILD_EXAMPLES OFF)
+endif()
+
 if(BUILD_SHARED_LIBS AND (WIN32 OR CYGWIN) AND (GINKGO_BUILD_TESTS OR GINKGO_BUILD_EXAMPLES OR GINKGO_BUILD_BENCHMARKS))
     # Change shared libraries output only if this build has executable program
     # with shared libraries.

benchmark/CMakeLists.txt

@@ -11,6 +11,12 @@ if (GINKGO_BUILD_CUDA AND GINKGO_BUILD_HIP AND GINKGO_HIP_PLATFORM MATCHES "hcc"
         "or use `export HIP_PLATFORM=nvcc` in your build environment instead.")
 endif()
 
+function(ginkgo_benchmark_add_tuning_maybe name)
+    if(GINKGO_BENCHMARK_ENABLE_TUNING)
+        target_sources(${name} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../utils/tuning_variables.cpp)
+    endif()
+endfunction()
+
 function(ginkgo_benchmark_cusp_linops name)
     target_compile_definitions("${name}" PRIVATE HAS_CUDA=1)
     target_link_libraries("${name}" ginkgo ${CUDA_RUNTIME_LIBS}
@@ -69,6 +75,7 @@ function(ginkgo_add_single_benchmark_executable name use_lib_linops macro_def)
     add_executable("${name}" ${ARGN})
     target_link_libraries("${name}" ginkgo gflags rapidjson)
     target_compile_definitions("${name}" PRIVATE "${macro_def}")
+    ginkgo_benchmark_add_tuning_maybe("${name}")
     if("${use_lib_linops}")
         if (GINKGO_BUILD_CUDA)
             ginkgo_benchmark_cusp_linops("${name}")

benchmark/conversions/conversions.cpp

@@ -51,6 +51,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "benchmark/utils/types.hpp"
 
 
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
+
 // This function supposes that management of `FLAGS_overwrite` is done before
 // calling it
 void convert_matrix(const gko::LinOp *matrix_from, const char *format_to,

benchmark/matrix_generator/matrix_generator.cpp

@@ -43,6 +43,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "benchmark/utils/types.hpp"
 
 
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
+
 namespace {
 std::string input_format =
     "  [\n"

benchmark/matrix_statistics/matrix_statistics.cpp

@@ -45,6 +45,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "benchmark/utils/types.hpp"
 
 
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
+
 // See en.wikipedia.org/wiki/Five-number_summary
 // Quartile computation uses Method 3 from en.wikipedia.org/wiki/Quartile
 void compute_summary(const std::vector<gko::size_type> &dist,

benchmark/preconditioner/preconditioner.cpp

@@ -51,8 +51,12 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "benchmark/utils/types.hpp"
 
 
-// preconditioner generation and application
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
 
+// preconditioner generation and application
 std::string encode_parameters(const char *precond_name)
 {
     static std::map<std::string, std::string (*)()> encoder{

benchmark/solver/solver.cpp

@@ -55,6 +55,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "benchmark/utils/types.hpp"
 
 
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
+
 // Command-line arguments
 DEFINE_uint32(max_iters, 1000,
               "Maximal number of iterations the solver will be run for");

benchmark/spmv/spmv.cpp

@@ -51,8 +51,12 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "benchmark/utils/types.hpp"
 
 
-// Command-line arguments
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
 
+// Command-line arguments
 DEFINE_uint32(nrhs, 1, "The number of right hand sides");
 
 
@@ -94,6 +98,47 @@ void apply_spmv(const char *format_name, std::shared_ptr<gko::Executor> exec,
             system_matrix->apply(lend(b), lend(x_clone));
             exec->synchronize();
         }
+
+        // tuning run
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+        auto &format_case = spmv_case[format_name];
+        if (!format_case.HasMember("tuning")) {
+            format_case.AddMember(
+                "tuning", rapidjson::Value(rapidjson::kObjectType), allocator);
+        }
+        auto &tuning_case = format_case["tuning"];
+        add_or_set_member(tuning_case, "time",
+                          rapidjson::Value(rapidjson::kArrayType), allocator);
+        add_or_set_member(tuning_case, "values",
+                          rapidjson::Value(rapidjson::kArrayType), allocator);
+
+        // Enable tuning for this portion of code
+        gko::_tuning_flag = true;
+        // Select some values we want to tune.
+        std::vector<gko::size_type> tuning_values{
+            1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096};
+        for (auto val : tuning_values) {
+            // Actually set the value that will be tuned. See
+            // cuda/components/format_conversion.cuh for an example of how this
+            // variable is used.
+            gko::_tuned_value = val;
+            auto tuning_timer = get_timer(exec, FLAGS_gpu_timer);
+            for (unsigned int i = 0; i < FLAGS_repetitions; i++) {
+                auto x_clone = clone(x);
+                exec->synchronize();
+                tuning_timer->tic();
+                system_matrix->apply(lend(b), lend(x_clone));
+                tuning_timer->toc();
+            }
+            tuning_case["time"].PushBack(tuning_timer->compute_average_time(),
+                                         allocator);
+            tuning_case["values"].PushBack(val, allocator);
+        }
+        // We put back the flag to false to use the default (non-tuned) values
+        // for the following
+        gko::_tuning_flag = false;
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
         // timed run
         auto timer = get_timer(exec, FLAGS_gpu_timer);
         for (unsigned int i = 0; i < FLAGS_repetitions; i++) {

benchmark/utils/tuning_variables.cpp

@@ -0,0 +1,48 @@
+/*******************************<GINKGO LICENSE>******************************
+Copyright (c) 2017-2021, the Ginkgo authors
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+******************************<GINKGO LICENSE>*******************************/
+
+#include <ginkgo/core/base/types.hpp>
+
+
+#include "benchmark/utils/tuning_variables.hpp"
+
+
+namespace gko {
+
+
+bool _tuning_flag = false;
+
+
+size_type _tuned_value = 0;
+
+
+}  // namespace gko

benchmark/utils/tuning_variables.hpp

@@ -0,0 +1,52 @@
+/*******************************<GINKGO LICENSE>******************************
+Copyright (c) 2017-2021, the Ginkgo authors
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+******************************<GINKGO LICENSE>*******************************/
+
+#ifndef GKO_BENCHMARK_UTILS_TUNING_VARIABLES_HPP_
+#define GKO_BENCHMARK_UTILS_TUNING_VARIABLES_HPP_
+
+
+#include <ginkgo/core/base/types.hpp>
+
+
+namespace gko {
+
+
+extern bool _tuning_flag;
+
+
+extern size_type _tuned_value;
+
+
+}  // namespace gko
+
+
+#endif  // GKO_BENCHMARK_UTILS_TUNING_VARIABLES_HPP_

cmake/get_info.cmake

@@ -130,7 +130,7 @@ foreach(log_type ${log_types})
         "GINKGO_BUILD_OMP;GINKGO_BUILD_REFERENCE;GINKGO_BUILD_CUDA;GINKGO_BUILD_HIP;GINKGO_BUILD_DPCPP")
     ginkgo_print_module_footer(${${log_type}} "  Tests, benchmarks and examples:")
     ginkgo_print_foreach_variable(
-        "GINKGO_BUILD_TESTS;GINKGO_BUILD_EXAMPLES;GINKGO_EXTLIB_EXAMPLE;GINKGO_BUILD_BENCHMARKS")
+        "GINKGO_BUILD_TESTS;GINKGO_BUILD_EXAMPLES;GINKGO_EXTLIB_EXAMPLE;GINKGO_BUILD_BENCHMARKS;GINKGO_BENCHMARK_ENABLE_TUNING")
     ginkgo_print_module_footer(${${log_type}} "  Documentation:")
     ginkgo_print_foreach_variable("GINKGO_BUILD_DOC;GINKGO_VERBOSE_LEVEL")
     ginkgo_print_module_footer(${${log_type}} "  Developer helpers:")

cuda/components/format_conversion.cuh

@@ -34,13 +34,19 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #define GKO_CUDA_COMPONENTS_FORMAT_CONVERSION_CUH_
 
 
+#include <ginkgo/config.hpp>
 #include <ginkgo/core/base/executor.hpp>
 
 
 #include "cuda/components/cooperative_groups.cuh"
 #include "cuda/components/thread_ids.cuh"
 
 
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
+
 namespace gko {
 namespace kernels {
 namespace cuda {
@@ -109,6 +115,11 @@ __host__ size_type calculate_nwarps(std::shared_ptr<const CudaExecutor> exec,
     } else if (nnz >= 2e5) {
         multiple = 32;
     }
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+    if (_tuning_flag) {
+        multiple = _tuned_value;
+    }
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
     return std::min(multiple * nwarps_in_cuda,
                     size_type(ceildiv(nnz, config::warp_size)));
 }

dev_tools/plots/tuning_heatmap.R

@@ -0,0 +1,49 @@
+# These packages are required, to install them, use the package manager or open
+# an R session and type:
+# install.packages("jsonlite", "tidyr", "ggplot2", "scales")
+library(jsonlite)
+library(ggplot2)
+library(scales)
+library(tidyr)
+
+# Manage arguments
+args <- commandArgs(trailingOnly=TRUE)
+if (length(args)!=2) {
+    stop("Usage: Rscript tuning_heatmap.R input_directory output_graphics_file\n", call=FALSE)
+}
+input <- args[1]
+output <- args[2]
+
+# Read the input json files into a dataframe
+files <- list.files(paste(input), recursive=TRUE, pattern = "*.json", full.names=TRUE)
+df_tmp <- list()
+count <- 1
+for (i in files)
+{
+    tmp <-jsonlite::fromJSON(i,flatten=TRUE)
+    df_tmp[[count]] <- as.data.frame(tmp)[,c("problem.name", "problem.nonzeros",
+                                             "spmv.coo.time", "spmv.coo.tuning.values",
+                                             "spmv.coo.tuning.time")]
+    count <- count +1
+}
+# Merge all the separate dataframes
+df_merged <- rbind_pages(df_tmp)
+# Unnest the two vectors
+df <- as.data.frame(unnest(df,spmv.coo.tuning.values,spmv.coo.tuning.time))
+# Now that all columns are vectors, compute the speedup using vector operations
+df$spmv.coo.speedup <- df$spmv.coo.time/df$spmv.coo.tuning.time
+
+# Plot the values
+ggplot(df, aes(factor(problem.nonzeros), factor(spmv.coo.tuning.values), fill=spmv.coo.speedup)) +
+    geom_tile() +
+    scale_fill_gradientn(
+        colours=c("red", "yellow", "skyblue", "darkblue"),
+        values = rescale(c(min(df$speedup),
+                           1.0,
+                           1.11,
+                           max(df$speedup)))) +
+    ggtitle("Speedup of tuned value against COO SpMV")+ xlab("nonzeros")+ ylab("tuned value (multiple)") +
+    theme(axis.text.x = element_text(angle = 90, vjust = 0.5, hjust=1), plot.title = element_text(hjust=0.5))
+
+# Save to the output file
+ggsave(paste(output), width=9, height=7)

hip/components/format_conversion.hip.hpp

@@ -37,13 +37,19 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <hip/hip_runtime.h>
 
 
+#include <ginkgo/config.hpp>
 #include <ginkgo/core/base/executor.hpp>
 
 
 #include "hip/components/cooperative_groups.hip.hpp"
 #include "hip/components/thread_ids.hip.hpp"
 
 
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+#include "benchmark/utils/tuning_variables.hpp"
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
+
+
 namespace gko {
 namespace kernels {
 namespace hip {
@@ -121,6 +127,11 @@ __host__ size_type calculate_nwarps(std::shared_ptr<const HipExecutor> exec,
         multiple = 8;
     }
 #endif  // GINKGO_HIP_PLATFORM_NVCC
+#ifdef GINKGO_BENCHMARK_ENABLE_TUNING
+    if (_tuning_flag) {
+        multiple = _tuned_value;
+    }
+#endif  // GINKGO_BENCHMARK_ENABLE_TUNING
     return std::min(multiple * nwarps_in_hip,
                     size_type(ceildiv(nnz, config::warp_size)));
 }