Tune and benchmark variables non intrusively.

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)));
 }