diff --git a/tile_engine/CMakeLists.txt b/tile_engine/CMakeLists.txt index b9dc320128..0bb885bc35 100644 --- a/tile_engine/CMakeLists.txt +++ b/tile_engine/CMakeLists.txt @@ -3,6 +3,7 @@ include_directories(BEFORE ${CMAKE_CURRENT_LIST_DIR}/include + ${CMAKE_CURRENT_LIST_DIR}/ops ) add_subdirectory(ops/gemm) diff --git a/tile_engine/ops/commons/__init__.py b/tile_engine/ops/common/__init__.py similarity index 100% rename from tile_engine/ops/commons/__init__.py rename to tile_engine/ops/common/__init__.py diff --git a/tile_engine/ops/commons/benchmark_utils.py b/tile_engine/ops/common/benchmark_utils.py similarity index 100% rename from tile_engine/ops/commons/benchmark_utils.py rename to tile_engine/ops/common/benchmark_utils.py diff --git a/tile_engine/ops/common/utils.hpp b/tile_engine/ops/common/utils.hpp new file mode 100644 index 0000000000..20994578e6 --- /dev/null +++ b/tile_engine/ops/common/utils.hpp @@ -0,0 +1,254 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once +#include +#include +#include +#include +#include +#include +#include + +#include "ck_tile/core.hpp" +#include "ck_tile/host.hpp" + +// Helper function to determine if a layout is row-major +template +constexpr auto is_row_major(Layout) +{ + return ck_tile::bool_constant>{}; +} + +// Structure to hold kernel traits for dispatcher +struct KernelTraits +{ + std::string pipeline; // compv3, compv4, mem + std::string scheduler; // intrawave, interwave + std::string epilogue; // cshuffle, default + bool pad_m; + bool pad_n; + bool pad_k; + bool persistent; + + // Constructor with defaults + KernelTraits() + : pipeline("compv3"), + scheduler("intrawave"), + epilogue("cshuffle"), + pad_m(false), + pad_n(false), + pad_k(false), + persistent(false) + { + } +}; + + +// Create argument parser +inline auto create_args(int argc, char* argv[]) +{ + ck_tile::ArgParser arg_parser; + arg_parser.insert("m", "3840", "The value for m dimension. Default is 3840.") + .insert("n", "4096", "The value for n dimension. Default is 4096.") + .insert("k", "2048", "The value for k dimension. Default is 2048.") + .insert("stride_a", "0", "The stride value for tensor A. Default is 0.") + .insert("stride_b", "0", "The stride value for tensor B. Default is 0.") + .insert("stride_ds", "0", "The stride value for tensor Ds . Default is 0.") + .insert("stride_c", "0", "The stride value for tensor C. Default is 0.") + .insert("split_k", "1", "The split value for k dimension. Default is 1.") + .insert("verify", + "2", + "The type of validation. Set to 0 for no validation, 1 for validation on CPU, or 2 " + "for validation on GPU. Default is 2, GPU validation.") + .insert("log", + "false", + "Whether output kernel instance information or not. Possible values are true or " + "false. Default is false") + .insert( + "warmup", "50", "The number of iterations before benchmark the kernel. Default is 50.") + .insert( + "repeat", "100", "The number of iterations to benchmark the kernel. Default is 100.") + .insert("timer", + "true", + "Whether if the timer is gpu timer or not. Possible values are false or true. " + "Default is true.") + .insert("init", + "0", + "The method of tensor initialization. Set to 0 for random, to 1 for linear, or 2 " + "for constant(1). Default is 0, random.") + .insert("flush_cache", + "true", + "To flush cache, possible values are true or false. " + "Default is false.") + .insert("rotating_count", "1000", "number of iterations to rotate the cache. default is 5.") + .insert("metric", + "0", + "Metric with which to measure kernel performance. Set to 0 for latency, 1 for " + "tflops, or 2 for bandwidth. Default is 0, latency.") + .insert("csv_filename", + "", + "The filename of benchmark result. Default is empty (no CSV output).") + .insert("structured_sparsity", + "false", + "Whether use sparsity kernel or not. Possible values are true or false. Default is " + "false") + .insert("json_output", + "false", + "Whether to output results in JSON format only. Possible values are true or false. " + "Default is " + "false"); + + bool result = arg_parser.parse(argc, argv); + return std::make_tuple(result, arg_parser); +} + +enum class Metric +{ + LATENCY = 0, + TFLOPS = 1, + BANDWIDTH = 2 +}; + +inline constexpr auto get_metric_name(Metric m) +{ + switch(m) + { + case Metric::LATENCY: return "latency"; + case Metric::TFLOPS: return "tflops"; + case Metric::BANDWIDTH: return "bandwidth"; + default: throw std::invalid_argument("Unsupported metric type"); + } +} + +struct PerformanceResult +{ + double latency_; + double tflops_; + double bandwidth_; + + static bool compare(const PerformanceResult& a, const PerformanceResult& b, Metric m) + { + switch(m) + { + case Metric::LATENCY: return a.latency_ < b.latency_; + case Metric::TFLOPS: return a.tflops_ > b.tflops_; + case Metric::BANDWIDTH: return a.bandwidth_ > b.bandwidth_; + default: throw std::invalid_argument("Unsupported metric type"); + } + } + + friend std::ostream& operator<<(std::ostream& os, const PerformanceResult& result) + { + os << "{\n" + << " \"latency(ms)\": " << std::fixed << std::setprecision(2) << result.latency_ + << ",\n" + << " \"tflops(TFlops)\": " << result.tflops_ << ",\n" + << " \"bandwidth(GB/s)\": " << result.bandwidth_ << "\n" + << "}"; + return os; + } +}; + +template +struct KernelInstance +{ + std::string name_; + Problem problem_; + PerformanceResult perf_result_; + + static bool compare(const KernelInstance& a, const KernelInstance& b, Metric m) + { + return PerformanceResult::compare(a.perf_result_, b.perf_result_, m); + } + + friend std::ostream& operator<<(std::ostream& os, const KernelInstance& obj) + { + os << "{\n" + << " \"name\": \"" << obj.name_ << "\",\n" + << " \"problem\": " << obj.problem_ << ",\n" + << " \"perf_result\": " << obj.perf_result_ << "\n" + << "}"; + return os; + } +}; + +struct Setting +{ + int n_warmup_; + int n_repeat_; + bool is_gpu_timer_; + int verify_; + int init_method_; + bool log_; + std::string csv_filename_; + bool flush_cache_; + int rotating_count_; + bool json_output_; +}; + +inline std::string get_rocm_version() +{ + std::ifstream version_file("/opt/rocm/.info/version"); + if(version_file.is_open()) + { + std::string version; + std::getline(version_file, version); + return version; + } + return "Unknown"; +} + +template +auto calculate_rtol_atol(const ck_tile::index_t K, + const ck_tile::index_t kbatch, + const float max_accumulated_value) +{ + using ComputeType = + std::conditional_t; + // Calculate thresholds + const auto rtol = ck_tile::get_relative_threshold( + ck_tile::integer_divide_ceil(K, kbatch)); + const auto atol = ck_tile::get_absolute_threshold( + max_accumulated_value / kbatch, ck_tile::integer_divide_ceil(K, kbatch)); + // Calculate error due to split_k accumulation + const auto rtol_split_k = + ck_tile::get_relative_threshold(kbatch); + const auto atol_split_k = ck_tile::get_absolute_threshold( + max_accumulated_value, kbatch); + // Use higher threshold + return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k)); +} + +template +auto calculate_rtol_atol(const ck_tile::index_t K, + const ck_tile::index_t kbatch, + const float max_accumulated_value) +{ + using ComputeTypeAB = + std::conditional_t; + + using ComputeType = + std::conditional_t; + + // Calculate thresholds + const auto rtol = ck_tile::get_relative_threshold( + ck_tile::integer_divide_ceil(K, kbatch)); + + const auto atol = ck_tile::get_absolute_threshold( + max_accumulated_value / kbatch, ck_tile::integer_divide_ceil(K, kbatch)); + + // Calculate error due to split_k accumulation + const auto rtol_split_k = + ck_tile::get_relative_threshold(kbatch); + + const auto atol_split_k = ck_tile::get_absolute_threshold( + max_accumulated_value, kbatch); + + // Use higher threshold + return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k)); +} diff --git a/tile_engine/ops/gemm/gemm_benchmark.hpp b/tile_engine/ops/gemm/gemm_benchmark.hpp new file mode 100644 index 0000000000..6ff09186a4 --- /dev/null +++ b/tile_engine/ops/gemm/gemm_benchmark.hpp @@ -0,0 +1,108 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include +#include +#include + +#include "ck_tile/core.hpp" +#include "ck_tile/host.hpp" +#include "common/utils.hpp" + +// Data types and Layouts are defined by the generated kernel headers +// No hardcoded type definitions here to avoid conflicts +struct GemmProblem +{ + int split_k_; + int m_, n_, k_; + int stride_a_, stride_b_, stride_c_; + + std::string dtype_a_, dtype_b_, dtype_acc_, dtype_c_; + std::string layout_a_, layout_b_, layout_c_; + + bool structured_sparsity_; + + friend std::ostream& operator<<(std::ostream& os, const GemmProblem& problem) + { + os << "{\n" + << " \"split_k\":" << problem.split_k_ << ",\n" + << " \"m\":" << problem.m_ << ",\n" + << " \"n\":" << problem.n_ << ",\n" + << " \"k\":" << problem.k_ << ",\n" + << " \"stride_a\":" << problem.stride_a_ << ",\n" + << " \"stride_b\":" << problem.stride_b_ << ",\n" + << " \"stride_c\":" << problem.stride_c_ << ",\n" + << " \"dtype_a\":\"" << problem.dtype_a_ << "\",\n" + << " \"dtype_b\":\"" << problem.dtype_b_ << "\",\n" + << " \"dtype_acc\":\"" << problem.dtype_acc_ << "\",\n" + << " \"dtype_c\":\"" << problem.dtype_c_ << "\",\n" + << " \"layout_a\":\"" << problem.layout_a_ << "\",\n" + << " \"layout_b\":\"" << problem.layout_b_ << "\",\n" + << " \"layout_c\":\"" << problem.layout_c_ << "\",\n" + << " \"structured_sparsity\":" << (problem.structured_sparsity_ ? "true" : "false") + << "\n" + << "}"; + return os; + } +}; + +// Detect Problem::DsDataType, default to void when absent +template +struct get_DsDataType { using type = void; }; + +template +struct get_DsDataType> { + using type = typename T::DsDataType; +}; + +// Detect Problem::D0DataType, default to void when absent +template +struct get_D0DataType { using type = void; }; + +template +struct get_D0DataType> { + using type = typename T::D0DataType; +}; + +/// @brief Function to compare the results of the device and host computations +template +bool compare(std::string instanceName, + ck_tile::index_t K, + ck_tile::index_t kbatch, + ck_tile::HostTensor& c_m_n_dev_result, + ck_tile::HostTensor& c_m_n_host_result) +{ + using DDataType = typename get_D0DataType::type; + const float max_accumulated_value = + *std::max_element(c_m_n_host_result.mData.begin(), c_m_n_host_result.mData.end()); + //const auto rtol_atol = calculate_rtol_atol( + //K, kbatch, max_accumulated_value); + auto rtol_atol = [&] { + if constexpr (std::is_void_v) + { + return calculate_rtol_atol( + K, kbatch, max_accumulated_value); + } + else + { + return calculate_rtol_atol( + K, kbatch, max_accumulated_value); + } + }(); + bool pass = ck_tile::check_err(c_m_n_dev_result, + c_m_n_host_result, + "Error: Incorrect results!", + rtol_atol.at(ck_tile::number<0>{}), + rtol_atol.at(ck_tile::number<1>{})); + + std::cout << "For " << instanceName << " Relative error threshold is " + << rtol_atol.at(ck_tile::number<0>{}) << " Absolute error threshold is " + << rtol_atol.at(ck_tile::number<1>{}) << std::endl; + std::cout << "The verification result is:" << (pass ? "correct" : "fail") << std::endl; + + return pass; +} diff --git a/tile_engine/ops/gemm/gemm_benchmark.py b/tile_engine/ops/gemm/gemm_benchmark.py index 3a7afc74e8..c229b14233 100644 --- a/tile_engine/ops/gemm/gemm_benchmark.py +++ b/tile_engine/ops/gemm/gemm_benchmark.py @@ -22,7 +22,7 @@ def _import_benchmark_utils(): # Load the module dynamically spec = importlib.util.spec_from_file_location( "benchmark_utils", - os.path.join(parent_dir, "commons", "benchmark_utils.py"), + os.path.join(parent_dir, "common", "benchmark_utils.py"), ) benchmark_utils = importlib.util.module_from_spec(spec) spec.loader.exec_module(benchmark_utils) diff --git a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp index f8c196e32a..23efbfc168 100644 --- a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp +++ b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp @@ -11,37 +11,15 @@ #include "ck_tile/core.hpp" #include "ck_tile/host.hpp" -#include "gemm_multi_d_common.hpp" +#include "gemm/gemm_benchmark.hpp" // Data types and Layouts are defined by the generated kernel headers // No hardcoded type definitions here to avoid conflicts - -enum class Metric +struct GemmMultiDProblem : GemmProblem { - LATENCY = 0, - TFLOPS = 1, - BANDWIDTH = 2 -}; - -inline constexpr auto get_metric_name(Metric m) -{ - switch(m) - { - case Metric::LATENCY: return "latency"; - case Metric::TFLOPS: return "tflops"; - case Metric::BANDWIDTH: return "bandwidth"; - default: throw std::invalid_argument("Unsupported metric type"); - } -} - -struct GemmMultiDProblem -{ - int split_k_; - int m_, n_, k_; - int stride_a_, stride_b_, stride_d0_, stride_d1_, stride_c_; - - std::string dtype_a_, dtype_b_, dtype_d0_, dtype_d1_, dtype_acc_, dtype_c_; - std::string layout_a_, layout_b_, layout_d0_, layout_d1_, layout_c_; + int stride_d0_, stride_d1_; + std::string dtype_d0_, dtype_d1_; + std::string layout_d0_, layout_d1_; friend std::ostream& operator<<(std::ostream& os, const GemmMultiDProblem& problem) { @@ -71,144 +49,6 @@ struct GemmMultiDProblem } }; -struct PerformanceResult -{ - double latency_; - double tflops_; - double bandwidth_; - - static bool compare(const PerformanceResult& a, const PerformanceResult& b, Metric m) - { - switch(m) - { - case Metric::LATENCY: return a.latency_ < b.latency_; - case Metric::TFLOPS: return a.tflops_ > b.tflops_; - case Metric::BANDWIDTH: return a.bandwidth_ > b.bandwidth_; - default: throw std::invalid_argument("Unsupported metric type"); - } - } - - friend std::ostream& operator<<(std::ostream& os, const PerformanceResult& result) - { - os << "{\n" - << " \"latency(ms)\": " << std::fixed << std::setprecision(2) << result.latency_ - << ",\n" - << " \"tflops(TFlops)\": " << result.tflops_ << ",\n" - << " \"bandwidth(GB/s)\": " << result.bandwidth_ << "\n" - << "}"; - return os; - } -}; - -struct KernelInstance -{ - std::string name_; - GemmMultiDProblem problem_; - PerformanceResult perf_result_; - - static bool compare(const KernelInstance& a, const KernelInstance& b, Metric m) - { - return PerformanceResult::compare(a.perf_result_, b.perf_result_, m); - } - - friend std::ostream& operator<<(std::ostream& os, const KernelInstance& obj) - { - os << "{\n" - << " \"name\": \"" << obj.name_ << "\",\n" - << " \"problem\": " << obj.problem_ << ",\n" - << " \"perf_result\": " << obj.perf_result_ << "\n" - << "}"; - return os; - } -}; - -struct Setting -{ - int n_warmup_; - int n_repeat_; - bool is_gpu_timer_; - int verify_; - int init_method_; - bool log_; - std::string csv_filename_; - bool flush_cache_; - int rotating_count_; - bool json_output_; -}; - -inline std::string get_rocm_version() -{ - std::ifstream version_file("/opt/rocm/.info/version"); - if(version_file.is_open()) - { - std::string version; - std::getline(version_file, version); - return version; - } - return "Unknown"; -} - -template -auto calculate_rtol_atol(const ck_tile::index_t K, - const ck_tile::index_t kbatch, - const float max_accumulated_value) -{ - using ComputeTypeAB = - std::conditional_t; - - using ComputeType = - std::conditional_t; - - // Calculate thresholds - const auto rtol = ck_tile::get_relative_threshold( - ck_tile::integer_divide_ceil(K, kbatch)); - - const auto atol = ck_tile::get_absolute_threshold( - max_accumulated_value / kbatch, ck_tile::integer_divide_ceil(K, kbatch)); - - // Calculate error due to split_k accumulation - const auto rtol_split_k = - ck_tile::get_relative_threshold(kbatch); - - const auto atol_split_k = ck_tile::get_absolute_threshold( - max_accumulated_value, kbatch); - - // Use higher threshold - return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k)); -} - -/// @brief Function to compare the results of the device and host computations -bool compare(std::string instanceName, - ck_tile::index_t K, - ck_tile::index_t kbatch, - ck_tile::HostTensor& c_m_n_dev_result, - ck_tile::HostTensor& c_m_n_host_result) -{ - const float max_accumulated_value = - *std::max_element(c_m_n_host_result.mData.begin(), c_m_n_host_result.mData.end()); - - const auto rtol_atol = - calculate_rtol_atol( - K, kbatch, max_accumulated_value); - - bool pass = ck_tile::check_err(c_m_n_dev_result, - c_m_n_host_result, - "Error: Incorrect results!", - rtol_atol.at(ck_tile::number<0>{}), - rtol_atol.at(ck_tile::number<1>{})); - - std::cout << "For " << instanceName << " Relative error threshold is " - << rtol_atol.at(ck_tile::number<0>{}) << " Absolute error threshold is " - << rtol_atol.at(ck_tile::number<1>{}) << std::endl; - std::cout << "The verification result is:" << (pass ? "correct" : "fail") << std::endl; - - return pass; -} - /// @brief Function to get the kernel output with reference implementation on CPU/GPU void gemm_multi_d_host_reference(int verify, ck_tile::HostTensor& a_m_k, diff --git a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.py b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.py index ae79668707..2e313c3fed 100644 --- a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.py +++ b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.py @@ -35,7 +35,7 @@ def _import_benchmark_utils(): # Load the module dynamically spec = importlib.util.spec_from_file_location( "benchmark_utils", - os.path.join(parent_dir, "commons", "benchmark_utils.py"), + os.path.join(parent_dir, "common", "benchmark_utils.py"), ) benchmark_utils = importlib.util.module_from_spec(spec) spec.loader.exec_module(benchmark_utils) diff --git a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark_single.cpp b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark_single.cpp index 41d2f736e1..e72ceb0d76 100644 --- a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark_single.cpp +++ b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark_single.cpp @@ -12,80 +12,20 @@ #include "ck_tile/core.hpp" #include "ck_tile/host.hpp" #include "gemm_multi_d_profiler.hpp" -#include "gemm_multi_d_common.hpp" // The kernel header is included via the compile command line with -include flag // It defines SelectedKernel struct and KERNEL_NAME -// DataTypeTraits are now defined in gemm_multi_d_common.hpp - -// Create argument parser -inline auto create_args(int argc, char* argv[]) -{ - ck_tile::ArgParser arg_parser; - arg_parser.insert("m", "3840", "The value for m dimension. Default is 3840.") - .insert("n", "4096", "The value for n dimension. Default is 4096.") - .insert("k", "2048", "The value for k dimension. Default is 2048.") - .insert("stride_a", "0", "The stride value for tensor A. Default is 0.") - .insert("stride_b", "0", "The stride value for tensor B. Default is 0.") - .insert("stride_ds", "0", "The stride value for tensor Ds . Default is 0.") - .insert("stride_c", "0", "The stride value for tensor C. Default is 0.") - .insert("split_k", "1", "The split value for k dimension. Default is 1.") - .insert("verify", - "1", - "for validation on GPU. Default is 1, validation on CPU, as validation on GPU is " - "not supported.") - .insert("log", - "false", - "Whether output kernel instance information or not. Possible values are true or " - "false. Default is false") - .insert( - "warmup", "50", "The number of iterations before benchmark the kernel. Default is 50.") - .insert( - "repeat", "100", "The number of iterations to benchmark the kernel. Default is 100.") - .insert("timer", - "true", - "Whether if the timer is gpu timer or not. Possible values are false or true. " - "Default is true.") - .insert("init", - "0", - "The method of tensor initialization. Set to 0 for random, to 1 for linear, or 2 " - "for constant(1). Default is 0, random.") - .insert("flush_cache", - "true", - "To flush cache, possible values are true or false. " - "Default is false.") - .insert("rotating_count", "1000", "number of iterations to rotate the cache. default is 5.") - .insert("metric", - "0", - "Metric with which to measure kernel performance. Set to 0 for latency, 1 for " - "tflops, or 2 for bandwidth. Default is 0, latency.") - .insert("csv_filename", - "", - "The filename of benchmark result. Default is empty (no CSV output).") - .insert("structured_sparsity", - "false", - "Whether use sparsity kernel or not. Possible values are true or false. Default is " - "false") - .insert("json_output", - "false", - "Whether to output results in JSON format only. Possible values are true or false. " - "Default is " - "false"); - - bool result = arg_parser.parse(argc, argv); - return std::make_tuple(result, arg_parser); -} void benchmark_single(const ck_tile::ArgParser& arg_parser) { // Use DataTypeTraits to get the actual type names from the generated header // The generated header defines ADataType, BDataType, AccDataType, CDataType - std::string dtype_a = DataTypeTraits::name; - std::string dtype_b = DataTypeTraits::name; - std::string dtype_acc = DataTypeTraits::name; - std::string dtype_c = DataTypeTraits::name; - std::string dtype_d0 = DataTypeTraits::name; - std::string dtype_d1 = DataTypeTraits::name; + std::string dtype_a = ck_tile::DataTypeTraits::name; + std::string dtype_b = ck_tile::DataTypeTraits::name; + std::string dtype_acc = ck_tile::DataTypeTraits::name; + std::string dtype_c = ck_tile::DataTypeTraits::name; + std::string dtype_d0 = ck_tile::DataTypeTraits::name; + std::string dtype_d1 = ck_tile::DataTypeTraits::name; // Layout names from the layout types std::string layout_a = ALayout::name; @@ -95,26 +35,29 @@ void benchmark_single(const ck_tile::ArgParser& arg_parser) std::string layout_d1 = D1Layout::name; // Create GemmMultiDProblem struct - GemmMultiDProblem gemm_multi_d_problem{arg_parser.get_int("split_k"), + GemmMultiDProblem gemm_multi_d_problem{ + GemmProblem{ + arg_parser.get_int("split_k"), arg_parser.get_int("m"), arg_parser.get_int("n"), arg_parser.get_int("k"), arg_parser.get_int("stride_a"), arg_parser.get_int("stride_b"), - arg_parser.get_int("stride_ds"), - arg_parser.get_int("stride_ds"), arg_parser.get_int("stride_c"), dtype_a, dtype_b, - dtype_d0, - dtype_d1, dtype_acc, dtype_c, layout_a, layout_b, - layout_d0, - layout_d1, - layout_c}; + layout_c, + arg_parser.get_bool("structured_sparsity")}, + arg_parser.get_int("stride_ds"), + arg_parser.get_int("stride_ds"), + dtype_d0, + dtype_d1, + layout_d0, + layout_d1}; // Create Setting struct Setting setting{arg_parser.get_int("warmup"), diff --git a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_common.hpp b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_common.hpp deleted file mode 100644 index 899221547f..0000000000 --- a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_common.hpp +++ /dev/null @@ -1,100 +0,0 @@ -// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT - -#pragma once - -#include -#include "ck_tile/core.hpp" -#include "ck_tile/host.hpp" -#include "ck_tile/core/numeric/integer.hpp" -#include "ck_tile/core/numeric/pk_int4.hpp" - -//[TODO] This can be moved to commons -// DataTypeTraits for all supported types -template -struct DataTypeTraits; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp32"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp64"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp16"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "bf16"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "bf8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "int8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "int32"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "pk_int4_t"; -}; - -// Helper function to determine if a layout is row-major -template -constexpr auto is_row_major(Layout) -{ - return ck_tile::bool_constant>{}; -} - -// Structure to hold kernel traits for dispatcher -struct KernelTraits -{ - std::string pipeline; // compv3, compv4, mem - std::string scheduler; // intrawave, interwave - std::string epilogue; // cshuffle, default - bool pad_m; - bool pad_n; - bool pad_k; - bool persistent; - - // Constructor with defaults - KernelTraits() - : pipeline("compv3"), - scheduler("intrawave"), - epilogue("cshuffle"), - pad_m(false), - pad_n(false), - pad_k(false), - persistent(false) - { - } -}; diff --git a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_profiler.hpp b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_profiler.hpp index 3a2cdc71fe..583dfd85a7 100644 --- a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_profiler.hpp +++ b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_profiler.hpp @@ -6,44 +6,36 @@ #include #include #include +#include +#include +#include +#include + #include "ck_tile/host/device_prop.hpp" #include "ck_tile/ops/gemm.hpp" +#include "gemm/gemm_profiler.hpp" #include "gemm_multi_d_benchmark.hpp" -class GemmMultiDProfiler +class GemmMultiDProfiler: public GemmProfiler> { - public: - static GemmMultiDProfiler& instance(Setting setting) - { - static GemmMultiDProfiler instance{setting}; - return instance; - } +public: + using BaseGemm = GemmProfiler>; + using BaseGemm::benchmark; - // Overload for single kernel benchmarking - void benchmark(GemmMultiDProblem& gemm_multi_d_problem, - std::function&, - const ck_tile::stream_config&)> kernel_func) - { - // Create a vector with a single callable that returns both name and time - std::vector( - ck_tile::GemmMultiDHostArgs&, const ck_tile::stream_config&)>> - callables; + GemmMultiDProfiler(Setting setting) + : GemmProfiler>(setting) {} - callables.push_back([kernel_func](ck_tile::GemmMultiDHostArgs& args, - const ck_tile::stream_config& stream) { - float time = kernel_func(args, stream); - return std::make_tuple(std::string(KERNEL_NAME), time); - }); - - benchmark(gemm_multi_d_problem, callables); - } void benchmark( GemmMultiDProblem& gemm_multi_d_problem, std::vector( ck_tile::GemmMultiDHostArgs&, const ck_tile::stream_config&)>>& - callables) + callables) override { const ALayout layout_a = ALayout{}; const BLayout layout_b = BLayout{}; @@ -166,142 +158,4 @@ class GemmMultiDProfiler } } - void process_result(const GemmMultiDProblem& gemm_multi_d_problem, - ck_tile::DeviceMem& c_m_n_dev_buf, - ck_tile::HostTensor& c_m_n_host_result, - ck_tile::HostTensor& c_m_n_dev_result, - const std::tuple& kernel_run_result) - { - auto [name, avg_time] = kernel_run_result; - - KernelInstance kernel_instance{name, gemm_multi_d_problem, {-1.0f, -1.0f, -1.0f}}; - - // compute performance metric - std::size_t flop = std::size_t(2) * gemm_multi_d_problem.m_ * gemm_multi_d_problem.n_ * - gemm_multi_d_problem.k_; - std::size_t num_byte = - sizeof(ADataType) * gemm_multi_d_problem.m_ * gemm_multi_d_problem.k_ + - sizeof(BDataType) * gemm_multi_d_problem.n_ * gemm_multi_d_problem.k_ + - sizeof(CDataType) * gemm_multi_d_problem.m_ * gemm_multi_d_problem.n_; - - // Dth Dimension Updates - ck_tile::static_for<0, DsDataType::size(), 1>{}([&](auto i) { - num_byte += sizeof(ck_tile::remove_cvref_t>) * - gemm_multi_d_problem.m_ * gemm_multi_d_problem.n_; - flop += sizeof(ck_tile::remove_cvref_t>) * - gemm_multi_d_problem.m_ * gemm_multi_d_problem.n_; - }); - - // update - kernel_instance.perf_result_.latency_ = avg_time; - kernel_instance.perf_result_.tflops_ = static_cast(flop) / 1.E9 / avg_time; - kernel_instance.perf_result_.bandwidth_ = num_byte / 1.E6 / avg_time; - - if(setting_.log_ > 0 && !setting_.json_output_) - { - std::cout << kernel_instance << std::endl; - } - - // verify result - c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data()); - bool verified_correct = - !setting_.verify_ || compare(name, - gemm_multi_d_problem.k_, - 1, // Multi d currently supports only k_batch = 1 - c_m_n_dev_result, - c_m_n_host_result); - - if(verified_correct) - { - kernel_instances_.emplace_back(kernel_instance); - } - else - { - std::cout << "Verification failed, skip kernel: " << name << std::endl; - } - - // clear tensor - c_m_n_dev_buf.SetZero(); - c_m_n_dev_result.SetZero(); - } - - KernelInstance select_best_instance(Metric metric) - { - if(kernel_instances_.empty()) - throw std::runtime_error("Empty instances"); - - auto kernel_instance = *std::max_element(kernel_instances_.begin(), - kernel_instances_.end(), - [metric](const auto& a, const auto& b) { - return PerformanceResult::compare( - b.perf_result_, a.perf_result_, metric); - }); - - if(setting_.json_output_) - { - // Output clean JSON only - std::cout << kernel_instance << std::endl; - } - else - { - std::cout << "**********************************" << std::endl; - std::cout << "According to given metrics: " << get_metric_name(metric) << "\n" - << "Current kernel performance is: " << kernel_instance << std::endl; - std::cout << "**********************************" << std::endl; - } - - if(!setting_.csv_filename_.empty()) - { - std::ofstream file(setting_.csv_filename_ + ".csv", std::ios::app); - - if(!file.is_open()) - { - std::cerr << "Warning: Failed to open CSV file for writing." << std::endl; - } - else - { - if(file.tellp() == 0) - { - file << "rocm_version,device_name," - << "split_k,m,n,k,stride_a,stride_b,stride_c," - << "dtype_a,dtype_b,dtype_acc,dtype_c," << "layout_a,layout_b,layout_c," - << "structured_sparsity," << "name," - << "latency(ms),tflops(TFlops),bandwidth(GB/s),metric\n"; - } - - const auto& problem = kernel_instance.problem_; - const auto& name = kernel_instance.name_; - const auto& perf = kernel_instance.perf_result_; - - file << get_rocm_version() << "," << ck_tile::get_device_name() << "," - << problem.split_k_ << "," << problem.m_ << "," << problem.n_ << "," - << problem.k_ << "," << problem.stride_a_ << "," << problem.stride_b_ << "," - << problem.stride_c_ << "," << problem.dtype_a_ << "," << problem.dtype_b_ - << "," << problem.dtype_acc_ << "," << problem.dtype_c_ << "," - << problem.layout_a_ << "," << problem.layout_b_ << "," << problem.layout_c_ - << "," << name << "," << std::fixed << std::setprecision(4) << perf.latency_ - << "," << std::fixed << std::setprecision(4) << perf.tflops_ << "," - << std::fixed << std::setprecision(4) << perf.bandwidth_ << "," - << get_metric_name(metric) << "\n"; - - if(!file) - { - std::cerr << "Warning: Error occurred while writing to CSV file." << std::endl; - } - } - } - - return kernel_instance; - } - - GemmMultiDProfiler(const GemmMultiDProfiler&) = delete; - GemmMultiDProfiler& operator=(const GemmMultiDProfiler&) = delete; - - private: - ~GemmMultiDProfiler() { kernel_instances_.clear(); } - GemmMultiDProfiler(Setting setting) : setting_(setting) {} - - Setting setting_; - - std::vector kernel_instances_; }; diff --git a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp index 748fe581d3..48f8a46ecb 100644 --- a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp +++ b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp @@ -6,191 +6,7 @@ #include "ck_tile/core.hpp" #include "ck_tile/host.hpp" #include "gemm_preshuffle_common.hpp" - -//[TODO] Move parts of this File to commons -enum class Metric -{ - LATENCY = 0, - TFLOPS = 1, - BANDWIDTH = 2 -}; - -inline constexpr auto get_metric_name(Metric m) -{ - switch(m) - { - case Metric::LATENCY: return "latency"; - case Metric::TFLOPS: return "tflops"; - case Metric::BANDWIDTH: return "bandwidth"; - default: throw std::invalid_argument("Unsupported metric type"); - } -} - -struct KernelConfig -{ - std::tuple tile_dims; - std::tuple warp_dims; - std::tuple warp_tile_dims; - bool permuteN; -}; - -struct GemmProblem -{ - int split_k_; - int m_, n_, k_; - int stride_a_, stride_b_, stride_c_; - - std::string dtype_a_, dtype_b_, dtype_acc_, dtype_c_; - std::string layout_a_, layout_b_, layout_c_; - - bool structured_sparsity_; - - friend std::ostream& operator<<(std::ostream& os, const GemmProblem& problem) - { - os << "{\n" - << " \"split_k\":" << problem.split_k_ << ",\n" - << " \"m\":" << problem.m_ << ",\n" - << " \"n\":" << problem.n_ << ",\n" - << " \"k\":" << problem.k_ << ",\n" - << " \"stride_a\":" << problem.stride_a_ << ",\n" - << " \"stride_b\":" << problem.stride_b_ << ",\n" - << " \"stride_c\":" << problem.stride_c_ << ",\n" - << " \"dtype_a\":\"" << problem.dtype_a_ << "\",\n" - << " \"dtype_b\":\"" << problem.dtype_b_ << "\",\n" - << " \"dtype_acc\":\"" << problem.dtype_acc_ << "\",\n" - << " \"dtype_c\":\"" << problem.dtype_c_ << "\",\n" - << " \"layout_a\":\"" << problem.layout_a_ << "\",\n" - << " \"layout_b\":\"" << problem.layout_b_ << "\",\n" - << " \"layout_c\":\"" << problem.layout_c_ << "\",\n" - << " \"structured_sparsity\":" << (problem.structured_sparsity_ ? "true" : "false") - << "\n" - << "}"; - return os; - } -}; - -struct PerformanceResult -{ - double latency_; - double tflops_; - double bandwidth_; - - static bool compare(const PerformanceResult& a, const PerformanceResult& b, Metric m) - { - switch(m) - { - case Metric::LATENCY: return a.latency_ < b.latency_; - case Metric::TFLOPS: return a.tflops_ > b.tflops_; - case Metric::BANDWIDTH: return a.bandwidth_ > b.bandwidth_; - default: throw std::invalid_argument("Unsupported metric type"); - } - } - - friend std::ostream& operator<<(std::ostream& os, const PerformanceResult& result) - { - os << "{\n" - << " \"latency(ms)\": " << std::fixed << std::setprecision(2) << result.latency_ - << ",\n" - << " \"tflops(TFlops)\": " << result.tflops_ << ",\n" - << " \"bandwidth(GB/s)\": " << result.bandwidth_ << "\n" - << "}"; - return os; - } -}; - -struct KernelInstance -{ - std::string name_; - GemmProblem problem_; - PerformanceResult perf_result_; - - static bool compare(const KernelInstance& a, const KernelInstance& b, Metric m) - { - return PerformanceResult::compare(a.perf_result_, b.perf_result_, m); - } - - friend std::ostream& operator<<(std::ostream& os, const KernelInstance& obj) - { - os << "{\n" - << " \"name\": \"" << obj.name_ << "\",\n" - << " \"problem\": " << obj.problem_ << ",\n" - << " \"perf_result\": " << obj.perf_result_ << "\n" - << "}"; - return os; - } -}; - -struct Setting -{ - int n_warmup_; - int n_repeat_; - bool is_gpu_timer_; - int verify_; - int init_method_; - bool log_; - std::string csv_filename_; - bool flush_cache_; - int rotating_count_; - bool json_output_; -}; - -inline std::string get_rocm_version() -{ - std::ifstream version_file("/opt/rocm/.info/version"); - if(version_file.is_open()) - { - std::string version; - std::getline(version_file, version); - return version; - } - return "Unknown"; -} - -template -auto calculate_rtol_atol(const ck_tile::index_t K, - const ck_tile::index_t kbatch, - const float max_accumulated_value) -{ - using ComputeType = - std::conditional_t; - // Calculate thresholds - const auto rtol = ck_tile::get_relative_threshold( - ck_tile::integer_divide_ceil(K, kbatch)); - const auto atol = ck_tile::get_absolute_threshold( - max_accumulated_value / kbatch, ck_tile::integer_divide_ceil(K, kbatch)); - // Calculate error due to split_k accumulation - const auto rtol_split_k = - ck_tile::get_relative_threshold(kbatch); - const auto atol_split_k = ck_tile::get_absolute_threshold( - max_accumulated_value, kbatch); - // Use higher threshold - return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k)); -} - -/// @brief Function to compare the results of the device and host computations -bool compare(std::string instanceName, - ck_tile::index_t K, - ck_tile::index_t kbatch, - ck_tile::HostTensor& c_m_n_dev_result, - ck_tile::HostTensor& c_m_n_ref) -{ - const float max_accumulated_value = - *std::max_element(c_m_n_ref.mData.begin(), c_m_n_ref.mData.end()); - const auto rtol_atol = calculate_rtol_atol( - K, kbatch, max_accumulated_value); - bool pass = ck_tile::check_err(c_m_n_dev_result, - c_m_n_ref, - "Error: Incorrect results!", - rtol_atol.at(ck_tile::number<0>{}), - rtol_atol.at(ck_tile::number<1>{})); - - std::cout << "For " << instanceName << " Relative error threshold is " - << rtol_atol.at(ck_tile::number<0>{}) << " Absolute error threshold is " - << rtol_atol.at(ck_tile::number<1>{}) << std::endl; - std::cout << "The verification result is:" << (pass ? "correct" : "fail") << std::endl; - - return pass; -} +#include "gemm/gemm_benchmark.hpp" /// @brief Function to get the kernel output with reference implementation on CPU/GPU void gemm_host_reference(int verify, diff --git a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.py b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.py index 1ea33834d7..935de186d6 100644 --- a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.py +++ b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.py @@ -36,7 +36,7 @@ def _import_benchmark_utils(): # Load the module dynamically spec = importlib.util.spec_from_file_location( "benchmark_utils", - os.path.join(parent_dir, "commons", "benchmark_utils.py"), + os.path.join(parent_dir, "common", "benchmark_utils.py"), ) benchmark_utils = importlib.util.module_from_spec(spec) spec.loader.exec_module(benchmark_utils) diff --git a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_common.hpp b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_common.hpp index 1b2cfe3735..7d3164f9d4 100644 --- a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_common.hpp +++ b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_common.hpp @@ -9,72 +9,6 @@ #include "ck_tile/core/numeric/integer.hpp" #include "ck_tile/core/numeric/pk_int4.hpp" -//[TODO] This can be moved to commons -// DataTypeTraits for all supported types -template -struct DataTypeTraits; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp32"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp64"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp16"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "bf16"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "bf8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "int8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "int32"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "pk_int4_t"; -}; - -// Helper function to determine if a layout is row-major -template -constexpr auto is_row_major(Layout) -{ - return ck_tile::bool_constant>{}; -} - // Structure to hold kernel traits for dispatcher struct KernelTraits { diff --git a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_profiler.hpp b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_profiler.hpp index 739bd7e677..4cf980dbf7 100644 --- a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_profiler.hpp +++ b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_profiler.hpp @@ -5,41 +5,27 @@ #include "ck_tile/host/device_prop.hpp" #include "ck_tile/ops/gemm.hpp" +#include "gemm/gemm_profiler.hpp" #include "gemm_preshuffle_benchmark.hpp" -class GemmProfiler +class GemmPreshuffleProfiler : public GemmProfiler { - public: - static GemmProfiler& instance(Setting setting) - { - static GemmProfiler instance{setting}; - return instance; - } +public: + using BaseGemm = GemmProfiler; + using BaseGemm::benchmark; - // Overload for single kernel benchmarking - void benchmark(GemmProblem& gemm_problem, - std::function - kernel_func, - KernelConfig& config) - { - // Create a vector with a single callable that returns both name and time - std::vector(ck_tile::GemmHostArgs&, - const ck_tile::stream_config&)>> - callables; + GemmPreshuffleProfiler(Setting setting) + : GemmProfiler(setting) {} - callables.push_back( - [kernel_func](ck_tile::GemmHostArgs& args, const ck_tile::stream_config& stream) { - float time = kernel_func(args, stream); - return std::make_tuple(std::string(KERNEL_NAME), time); - }); - - benchmark(gemm_problem, callables, config); - } void benchmark(GemmProblem& gemm_problem, std::vector( ck_tile::GemmHostArgs&, const ck_tile::stream_config&)>>& callables, - KernelConfig& config) + KernelConfig& config) override { const ALayout layout_a = ALayout{}; const BLayout layout_b = BLayout{}; @@ -159,131 +145,4 @@ class GemmProfiler } } - void process_result(const GemmProblem& gemm_problem, - ck_tile::DeviceMem& c_m_n_dev_buf, - ck_tile::HostTensor& c_m_n_ref, - ck_tile::HostTensor& c_m_n_dev_result, - const std::tuple& kernel_run_result) - { - auto [name, avg_time] = kernel_run_result; - - KernelInstance kernel_instance{name, gemm_problem, {-1.0f, -1.0f, -1.0f}}; - - // compute performance metric - std::size_t flop = std::size_t(2) * gemm_problem.m_ * gemm_problem.n_ * gemm_problem.k_; - std::size_t num_byte = sizeof(ADataType) * gemm_problem.m_ * gemm_problem.k_ + - sizeof(BDataType) * gemm_problem.n_ * gemm_problem.k_ + - sizeof(CDataType) * gemm_problem.m_ * gemm_problem.n_; - - // update - kernel_instance.perf_result_.latency_ = avg_time; - kernel_instance.perf_result_.tflops_ = static_cast(flop) / 1.E9 / avg_time; - kernel_instance.perf_result_.bandwidth_ = num_byte / 1.E6 / avg_time; - - if(setting_.log_ > 0 && !setting_.json_output_) - { - std::cout << kernel_instance << std::endl; - } - - // verify result - c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data()); - - bool verified_correct = - !setting_.verify_ || - compare(name, gemm_problem.k_, gemm_problem.split_k_, c_m_n_dev_result, c_m_n_ref); - - if(verified_correct) - { - kernel_instances_.emplace_back(kernel_instance); - } - else - { - std::cout << "Verification failed, skip kernel: " << name << std::endl; - } - - // clear tensor - c_m_n_dev_buf.SetZero(); - c_m_n_dev_result.SetZero(); - } - - KernelInstance select_best_instance(Metric metric) - { - if(kernel_instances_.empty()) - throw std::runtime_error("Empty instances"); - - auto kernel_instance = *std::max_element(kernel_instances_.begin(), - kernel_instances_.end(), - [metric](const auto& a, const auto& b) { - return PerformanceResult::compare( - b.perf_result_, a.perf_result_, metric); - }); - - if(setting_.json_output_) - { - // Output clean JSON only - std::cout << kernel_instance << std::endl; - } - else - { - std::cout << "**********************************" << std::endl; - std::cout << "According to given metrics: " << get_metric_name(metric) << "\n" - << "Current kernel performance is: " << kernel_instance << std::endl; - std::cout << "**********************************" << std::endl; - } - - if(!setting_.csv_filename_.empty()) - { - std::ofstream file(setting_.csv_filename_ + ".csv", std::ios::app); - - if(!file.is_open()) - { - std::cerr << "Warning: Failed to open CSV file for writing." << std::endl; - } - else - { - if(file.tellp() == 0) - { - file << "rocm_version,device_name," - << "split_k,m,n,k,stride_a,stride_b,stride_c," - << "dtype_a,dtype_b,dtype_acc,dtype_c," << "layout_a,layout_b,layout_c," - << "structured_sparsity," << "name," - << "latency(ms),tflops(TFlops),bandwidth(GB/s),metric\n"; - } - - const auto& problem = kernel_instance.problem_; - const auto& name = kernel_instance.name_; - const auto& perf = kernel_instance.perf_result_; - - file << get_rocm_version() << "," << ck_tile::get_device_name() << "," - << problem.split_k_ << "," << problem.m_ << "," << problem.n_ << "," - << problem.k_ << "," << problem.stride_a_ << "," << problem.stride_b_ << "," - << problem.stride_c_ << "," << problem.dtype_a_ << "," << problem.dtype_b_ - << "," << problem.dtype_acc_ << "," << problem.dtype_c_ << "," - << problem.layout_a_ << "," << problem.layout_b_ << "," << problem.layout_c_ - << "," << problem.structured_sparsity_ << "," << name << "," << std::fixed - << std::setprecision(4) << perf.latency_ << "," << std::fixed - << std::setprecision(4) << perf.tflops_ << "," << std::fixed - << std::setprecision(4) << perf.bandwidth_ << "," << get_metric_name(metric) - << "\n"; - - if(!file) - { - std::cerr << "Warning: Error occurred while writing to CSV file." << std::endl; - } - } - } - - return kernel_instance; - } - - GemmProfiler(const GemmProfiler&) = delete; - GemmProfiler& operator=(const GemmProfiler&) = delete; - - private: - ~GemmProfiler() { kernel_instances_.clear(); } - GemmProfiler(Setting setting) : setting_(setting) {} - - Setting setting_; - - std::vector kernel_instances_; }; diff --git a/tile_engine/ops/gemm/gemm_profiler.hpp b/tile_engine/ops/gemm/gemm_profiler.hpp new file mode 100644 index 0000000000..4c9b706fe3 --- /dev/null +++ b/tile_engine/ops/gemm/gemm_profiler.hpp @@ -0,0 +1,200 @@ +// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include +#include +#include +#include +#include + + +#include "ck_tile/host/device_prop.hpp" +#include "ck_tile/ops/gemm.hpp" +#include "gemm_benchmark.hpp" + +template +class GemmProfiler +{ +public: + static Gemm& instance(Setting setting) + { + static Gemm instance{setting}; + return instance; + } + + // Overload for single kernel benchmarking + void benchmark(Problem& gemm_problem, + std::function + kernel_func) + { + // Create a vector with a single callable that returns both name and time + std::vector(GemmArgs&, + const ck_tile::stream_config&)>> + callables; + + callables.push_back( + [kernel_func](GemmArgs& args, const ck_tile::stream_config& stream) { + float time = kernel_func(args, stream); + return std::make_tuple(std::string(KERNEL_NAME), time); + }); + + benchmark(gemm_problem, callables); // TODO: need to cast this? + } + + virtual void benchmark(Problem& gemm_problem, + std::vector( + GemmArgs&, const ck_tile::stream_config&)>>& callables) = 0; + + void process_result(const Problem& gemm_problem, + ck_tile::DeviceMem& c_m_n_dev_buf, + ck_tile::HostTensor& c_m_n_host_result, + ck_tile::HostTensor& c_m_n_dev_result, + const std::tuple& kernel_run_result) + { + auto [name, avg_time] = kernel_run_result; + using DDataType = typename get_DsDataType::type; + + KernelInstance kernel_instance{name, gemm_problem, {-1.0f, -1.0f, -1.0f}}; + + // compute performance metric + std::size_t flop = std::size_t(2) * gemm_problem.m_ * gemm_problem.n_ * gemm_problem.k_; + std::size_t num_byte = sizeof(ADataType) * gemm_problem.m_ * gemm_problem.k_ + + sizeof(BDataType) * gemm_problem.n_ * gemm_problem.k_ + + sizeof(CDataType) * gemm_problem.m_ * gemm_problem.n_; + + + if constexpr (!std::is_void_v) + { + ck_tile::static_for<0, DDataType::size(), 1>{}([&](auto i) { + using DType = ck_tile::remove_cvref_t>; + num_byte += sizeof(DType) * gemm_problem.m_ * gemm_problem.n_; + flop += sizeof(DType) * gemm_problem.m_ * gemm_problem.n_; + }); + } + + + // update + kernel_instance.perf_result_.latency_ = avg_time; + kernel_instance.perf_result_.tflops_ = static_cast(flop) / 1.E9 / avg_time; + kernel_instance.perf_result_.bandwidth_ = num_byte / 1.E6 / avg_time; + + if(setting_.log_ > 0 && !setting_.json_output_) + { + std::cout << kernel_instance << std::endl; + } + + // verify result + c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data()); + int split_k = 1; + if constexpr (std::is_same_v) + { + split_k = gemm_problem.split_k_; + } + bool verified_correct = + !setting_.verify_ || + compare( + name, gemm_problem.k_, split_k, c_m_n_dev_result, c_m_n_host_result); + + if(verified_correct) + { + kernel_instances_.emplace_back(kernel_instance); + } + else + { + std::cout << "Verification failed, skip kernel: " << name << std::endl; + } + + // clear tensor + c_m_n_dev_buf.SetZero(); + c_m_n_dev_result.SetZero(); + } + + KernelInstance select_best_instance(Metric metric) + { + if(kernel_instances_.empty()) + throw std::runtime_error("Empty instances"); + + auto kernel_instance = *std::max_element(kernel_instances_.begin(), + kernel_instances_.end(), + [metric](const auto& a, const auto& b) { + return PerformanceResult::compare( + b.perf_result_, a.perf_result_, metric); + }); + + if(setting_.json_output_) + { + // Output clean JSON only + std::cout << kernel_instance << std::endl; + } + else + { + std::cout << "**********************************" << std::endl; + std::cout << "According to given metrics: " << get_metric_name(metric) << "\n" + << "Current kernel performance is: " << kernel_instance << std::endl; + std::cout << "**********************************" << std::endl; + } + + if(!setting_.csv_filename_.empty()) + { + std::ofstream file(setting_.csv_filename_ + ".csv", std::ios::app); + + if(!file.is_open()) + { + std::cerr << "Warning: Failed to open CSV file for writing." << std::endl; + } + else + { + if(file.tellp() == 0) + { + file << "rocm_version,device_name," + << "split_k,m,n,k,stride_a,stride_b,stride_c," + << "dtype_a,dtype_b,dtype_acc,dtype_c," << "layout_a,layout_b,layout_c," + << "structured_sparsity," << "name," + << "latency(ms),tflops(TFlops),bandwidth(GB/s),metric\n"; + } + + const auto& problem = kernel_instance.problem_; + const auto& name = kernel_instance.name_; + const auto& perf = kernel_instance.perf_result_; + + file << get_rocm_version() << "," << ck_tile::get_device_name() << "," + << problem.split_k_ << "," << problem.m_ << "," << problem.n_ << "," + << problem.k_ << "," << problem.stride_a_ << "," << problem.stride_b_ << "," + << problem.stride_c_ << "," << problem.dtype_a_ << "," << problem.dtype_b_ + << "," << problem.dtype_acc_ << "," << problem.dtype_c_ << "," + << problem.layout_a_ << "," << problem.layout_b_ << "," << problem.layout_c_ + << "," << problem.structured_sparsity_ << "," << name << "," << std::fixed + << std::setprecision(4) << perf.latency_ << "," << std::fixed + << std::setprecision(4) << perf.tflops_ << "," << std::fixed + << std::setprecision(4) << perf.bandwidth_ << "," << get_metric_name(metric) + << "\n"; + + if(!file) + { + std::cerr << "Warning: Error occurred while writing to CSV file." << std::endl; + } + } + } + + return kernel_instance; + } + + GemmProfiler(const GemmProfiler&) = delete; + GemmProfiler& operator=(const GemmProfiler&) = delete; + + protected: + virtual ~GemmProfiler() { kernel_instances_.clear(); } + GemmProfiler(Setting setting) : setting_(setting) {} + + Setting setting_; + + std::vector> kernel_instances_; +}; + + diff --git a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp index c7f4f470b0..4cf272a3c8 100644 --- a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp +++ b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp @@ -11,187 +11,11 @@ #include "ck_tile/core.hpp" #include "ck_tile/host.hpp" -#include "gemm_universal_common.hpp" +#include "gemm/gemm_benchmark.hpp" // Data types and Layouts are defined by the generated kernel headers // No hardcoded type definitions here to avoid conflicts -enum class Metric -{ - LATENCY = 0, - TFLOPS = 1, - BANDWIDTH = 2 -}; - -inline constexpr auto get_metric_name(Metric m) -{ - switch(m) - { - case Metric::LATENCY: return "latency"; - case Metric::TFLOPS: return "tflops"; - case Metric::BANDWIDTH: return "bandwidth"; - default: throw std::invalid_argument("Unsupported metric type"); - } -} - -struct GemmProblem -{ - int split_k_; - int m_, n_, k_; - int stride_a_, stride_b_, stride_c_; - - std::string dtype_a_, dtype_b_, dtype_acc_, dtype_c_; - std::string layout_a_, layout_b_, layout_c_; - - bool structured_sparsity_; - - friend std::ostream& operator<<(std::ostream& os, const GemmProblem& problem) - { - os << "{\n" - << " \"split_k\":" << problem.split_k_ << ",\n" - << " \"m\":" << problem.m_ << ",\n" - << " \"n\":" << problem.n_ << ",\n" - << " \"k\":" << problem.k_ << ",\n" - << " \"stride_a\":" << problem.stride_a_ << ",\n" - << " \"stride_b\":" << problem.stride_b_ << ",\n" - << " \"stride_c\":" << problem.stride_c_ << ",\n" - << " \"dtype_a\":\"" << problem.dtype_a_ << "\",\n" - << " \"dtype_b\":\"" << problem.dtype_b_ << "\",\n" - << " \"dtype_acc\":\"" << problem.dtype_acc_ << "\",\n" - << " \"dtype_c\":\"" << problem.dtype_c_ << "\",\n" - << " \"layout_a\":\"" << problem.layout_a_ << "\",\n" - << " \"layout_b\":\"" << problem.layout_b_ << "\",\n" - << " \"layout_c\":\"" << problem.layout_c_ << "\",\n" - << " \"structured_sparsity\":" << (problem.structured_sparsity_ ? "true" : "false") - << "\n" - << "}"; - return os; - } -}; - -struct PerformanceResult -{ - double latency_; - double tflops_; - double bandwidth_; - - static bool compare(const PerformanceResult& a, const PerformanceResult& b, Metric m) - { - switch(m) - { - case Metric::LATENCY: return a.latency_ < b.latency_; - case Metric::TFLOPS: return a.tflops_ > b.tflops_; - case Metric::BANDWIDTH: return a.bandwidth_ > b.bandwidth_; - default: throw std::invalid_argument("Unsupported metric type"); - } - } - - friend std::ostream& operator<<(std::ostream& os, const PerformanceResult& result) - { - os << "{\n" - << " \"latency(ms)\": " << std::fixed << std::setprecision(2) << result.latency_ - << ",\n" - << " \"tflops(TFlops)\": " << result.tflops_ << ",\n" - << " \"bandwidth(GB/s)\": " << result.bandwidth_ << "\n" - << "}"; - return os; - } -}; - -struct KernelInstance -{ - std::string name_; - GemmProblem problem_; - PerformanceResult perf_result_; - - static bool compare(const KernelInstance& a, const KernelInstance& b, Metric m) - { - return PerformanceResult::compare(a.perf_result_, b.perf_result_, m); - } - - friend std::ostream& operator<<(std::ostream& os, const KernelInstance& obj) - { - os << "{\n" - << " \"name\": \"" << obj.name_ << "\",\n" - << " \"problem\": " << obj.problem_ << ",\n" - << " \"perf_result\": " << obj.perf_result_ << "\n" - << "}"; - return os; - } -}; - -struct Setting -{ - int n_warmup_; - int n_repeat_; - bool is_gpu_timer_; - int verify_; - int init_method_; - bool log_; - std::string csv_filename_; - bool flush_cache_; - int rotating_count_; - bool json_output_; -}; - -inline std::string get_rocm_version() -{ - std::ifstream version_file("/opt/rocm/.info/version"); - if(version_file.is_open()) - { - std::string version; - std::getline(version_file, version); - return version; - } - return "Unknown"; -} - -template -auto calculate_rtol_atol(const ck_tile::index_t K, - const ck_tile::index_t kbatch, - const float max_accumulated_value) -{ - using ComputeType = - std::conditional_t; - // Calculate thresholds - const auto rtol = ck_tile::get_relative_threshold( - ck_tile::integer_divide_ceil(K, kbatch)); - const auto atol = ck_tile::get_absolute_threshold( - max_accumulated_value / kbatch, ck_tile::integer_divide_ceil(K, kbatch)); - // Calculate error due to split_k accumulation - const auto rtol_split_k = - ck_tile::get_relative_threshold(kbatch); - const auto atol_split_k = ck_tile::get_absolute_threshold( - max_accumulated_value, kbatch); - // Use higher threshold - return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k)); -} - -/// @brief Function to compare the results of the device and host computations -bool compare(std::string instanceName, - ck_tile::index_t K, - ck_tile::index_t kbatch, - ck_tile::HostTensor& c_m_n_dev_result, - ck_tile::HostTensor& c_m_n_host_result) -{ - const float max_accumulated_value = - *std::max_element(c_m_n_host_result.mData.begin(), c_m_n_host_result.mData.end()); - const auto rtol_atol = calculate_rtol_atol( - K, kbatch, max_accumulated_value); - bool pass = ck_tile::check_err(c_m_n_dev_result, - c_m_n_host_result, - "Error: Incorrect results!", - rtol_atol.at(ck_tile::number<0>{}), - rtol_atol.at(ck_tile::number<1>{})); - - std::cout << "For " << instanceName << " Relative error threshold is " - << rtol_atol.at(ck_tile::number<0>{}) << " Absolute error threshold is " - << rtol_atol.at(ck_tile::number<1>{}) << std::endl; - std::cout << "The verification result is:" << (pass ? "correct" : "fail") << std::endl; - - return pass; -} - /// @brief Function to get the kernel output with reference implementation on CPU/GPU void gemm_host_reference(int verify, ck_tile::HostTensor& a_m_k, diff --git a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.py b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.py index 88ed4465af..4fa5d5dee0 100755 --- a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.py +++ b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.py @@ -36,7 +36,7 @@ def _import_benchmark_utils(): # Load the module dynamically spec = importlib.util.spec_from_file_location( "benchmark_utils", - os.path.join(parent_dir, "commons", "benchmark_utils.py"), + os.path.join(parent_dir, "common", "benchmark_utils.py"), ) benchmark_utils = importlib.util.module_from_spec(spec) spec.loader.exec_module(benchmark_utils) diff --git a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark_single.cpp b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark_single.cpp index 613a42ff80..bf81acabd2 100644 --- a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark_single.cpp +++ b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark_single.cpp @@ -12,20 +12,18 @@ #include "ck_tile/core.hpp" #include "ck_tile/host.hpp" #include "gemm_universal_profiler.hpp" -#include "gemm_universal_common.hpp" // The kernel header is included via the compile command line with -include flag // It defines SelectedKernel struct and KERNEL_NAME -// DataTypeTraits are now defined in gemm_common.hpp void benchmark_single(const ck_tile::ArgParser& arg_parser) { // Use DataTypeTraits to get the actual type names from the generated header // The generated header defines ADataType, BDataType, AccDataType, CDataType - std::string dtype_a = DataTypeTraits::name; - std::string dtype_b = DataTypeTraits::name; - std::string dtype_acc = DataTypeTraits::name; - std::string dtype_c = DataTypeTraits::name; + std::string dtype_a = ck_tile::DataTypeTraits::name; + std::string dtype_b = ck_tile::DataTypeTraits::name; + std::string dtype_acc = ck_tile::DataTypeTraits::name; + std::string dtype_c = ck_tile::DataTypeTraits::name; // Layout names from the layout types std::string layout_a = ALayout::name; @@ -62,7 +60,7 @@ void benchmark_single(const ck_tile::ArgParser& arg_parser) arg_parser.get_bool("json_output")}; // Get the profiler instance - auto& profiler = GemmProfiler::instance(setting); + auto& profiler = UniversalGemmProfiler::GemmProfiler::instance(setting); try { diff --git a/tile_engine/ops/gemm/gemm_universal/gemm_universal_common.hpp b/tile_engine/ops/gemm/gemm_universal/gemm_universal_common.hpp deleted file mode 100644 index 899221547f..0000000000 --- a/tile_engine/ops/gemm/gemm_universal/gemm_universal_common.hpp +++ /dev/null @@ -1,100 +0,0 @@ -// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. -// SPDX-License-Identifier: MIT - -#pragma once - -#include -#include "ck_tile/core.hpp" -#include "ck_tile/host.hpp" -#include "ck_tile/core/numeric/integer.hpp" -#include "ck_tile/core/numeric/pk_int4.hpp" - -//[TODO] This can be moved to commons -// DataTypeTraits for all supported types -template -struct DataTypeTraits; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp32"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp64"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp16"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "bf16"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "fp8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "bf8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "int8"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "int32"; -}; - -template <> -struct DataTypeTraits -{ - static constexpr const char* name = "pk_int4_t"; -}; - -// Helper function to determine if a layout is row-major -template -constexpr auto is_row_major(Layout) -{ - return ck_tile::bool_constant>{}; -} - -// Structure to hold kernel traits for dispatcher -struct KernelTraits -{ - std::string pipeline; // compv3, compv4, mem - std::string scheduler; // intrawave, interwave - std::string epilogue; // cshuffle, default - bool pad_m; - bool pad_n; - bool pad_k; - bool persistent; - - // Constructor with defaults - KernelTraits() - : pipeline("compv3"), - scheduler("intrawave"), - epilogue("cshuffle"), - pad_m(false), - pad_n(false), - pad_k(false), - persistent(false) - { - } -}; diff --git a/tile_engine/ops/gemm/gemm_universal/gemm_universal_profiler.hpp b/tile_engine/ops/gemm/gemm_universal/gemm_universal_profiler.hpp index 9b728c52d6..a1d749d678 100644 --- a/tile_engine/ops/gemm/gemm_universal/gemm_universal_profiler.hpp +++ b/tile_engine/ops/gemm/gemm_universal/gemm_universal_profiler.hpp @@ -10,38 +10,24 @@ #include "ck_tile/host/device_prop.hpp" #include "ck_tile/ops/gemm.hpp" #include "gemm_universal_benchmark.hpp" +#include "gemm/gemm_profiler.hpp" -class GemmProfiler +class UniversalGemmProfiler : public GemmProfiler { - public: - static GemmProfiler& instance(Setting setting) - { - static GemmProfiler instance{setting}; - return instance; - } +public: + using BaseGemm = GemmProfiler; + using BaseGemm::benchmark; - // Overload for single kernel benchmarking - void benchmark(GemmProblem& gemm_problem, - std::function - kernel_func) - { - // Create a vector with a single callable that returns both name and time - std::vector(ck_tile::GemmHostArgs&, - const ck_tile::stream_config&)>> - callables; - - callables.push_back( - [kernel_func](ck_tile::GemmHostArgs& args, const ck_tile::stream_config& stream) { - float time = kernel_func(args, stream); - return std::make_tuple(std::string(KERNEL_NAME), time); - }); - - benchmark(gemm_problem, callables); - } + UniversalGemmProfiler(Setting setting) + : GemmProfiler(setting) {} void benchmark(GemmProblem& gemm_problem, std::vector( - ck_tile::GemmHostArgs&, const ck_tile::stream_config&)>>& callables) + ck_tile::GemmHostArgs&, const ck_tile::stream_config&)>>& callables) override { const ALayout layout_a = ALayout{}; const BLayout layout_b = BLayout{}; @@ -158,132 +144,4 @@ class GemmProfiler kernel_run_result); } } - - void process_result(const GemmProblem& gemm_problem, - ck_tile::DeviceMem& c_m_n_dev_buf, - ck_tile::HostTensor& c_m_n_host_result, - ck_tile::HostTensor& c_m_n_dev_result, - const std::tuple& kernel_run_result) - { - auto [name, avg_time] = kernel_run_result; - - KernelInstance kernel_instance{name, gemm_problem, {-1.0f, -1.0f, -1.0f}}; - - // compute performance metric - std::size_t flop = std::size_t(2) * gemm_problem.m_ * gemm_problem.n_ * gemm_problem.k_; - std::size_t num_byte = sizeof(ADataType) * gemm_problem.m_ * gemm_problem.k_ + - sizeof(BDataType) * gemm_problem.n_ * gemm_problem.k_ + - sizeof(CDataType) * gemm_problem.m_ * gemm_problem.n_; - - // update - kernel_instance.perf_result_.latency_ = avg_time; - kernel_instance.perf_result_.tflops_ = static_cast(flop) / 1.E9 / avg_time; - kernel_instance.perf_result_.bandwidth_ = num_byte / 1.E6 / avg_time; - - if(setting_.log_ > 0 && !setting_.json_output_) - { - std::cout << kernel_instance << std::endl; - } - - // verify result - c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data()); - bool verified_correct = - !setting_.verify_ || - compare( - name, gemm_problem.k_, gemm_problem.split_k_, c_m_n_dev_result, c_m_n_host_result); - - if(verified_correct) - { - kernel_instances_.emplace_back(kernel_instance); - } - else - { - std::cout << "Verification failed, skip kernel: " << name << std::endl; - } - - // clear tensor - c_m_n_dev_buf.SetZero(); - c_m_n_dev_result.SetZero(); - } - - KernelInstance select_best_instance(Metric metric) - { - if(kernel_instances_.empty()) - throw std::runtime_error("Empty instances"); - - auto kernel_instance = *std::max_element(kernel_instances_.begin(), - kernel_instances_.end(), - [metric](const auto& a, const auto& b) { - return PerformanceResult::compare( - b.perf_result_, a.perf_result_, metric); - }); - - if(setting_.json_output_) - { - // Output clean JSON only - std::cout << kernel_instance << std::endl; - } - else - { - std::cout << "**********************************" << std::endl; - std::cout << "According to given metrics: " << get_metric_name(metric) << "\n" - << "Current kernel performance is: " << kernel_instance << std::endl; - std::cout << "**********************************" << std::endl; - } - - if(!setting_.csv_filename_.empty()) - { - std::ofstream file(setting_.csv_filename_ + ".csv", std::ios::app); - - if(!file.is_open()) - { - std::cerr << "Warning: Failed to open CSV file for writing." << std::endl; - } - else - { - if(file.tellp() == 0) - { - file << "rocm_version,device_name," - << "split_k,m,n,k,stride_a,stride_b,stride_c," - << "dtype_a,dtype_b,dtype_acc,dtype_c," << "layout_a,layout_b,layout_c," - << "structured_sparsity," << "name," - << "latency(ms),tflops(TFlops),bandwidth(GB/s),metric\n"; - } - - const auto& problem = kernel_instance.problem_; - const auto& name = kernel_instance.name_; - const auto& perf = kernel_instance.perf_result_; - - file << get_rocm_version() << "," << ck_tile::get_device_name() << "," - << problem.split_k_ << "," << problem.m_ << "," << problem.n_ << "," - << problem.k_ << "," << problem.stride_a_ << "," << problem.stride_b_ << "," - << problem.stride_c_ << "," << problem.dtype_a_ << "," << problem.dtype_b_ - << "," << problem.dtype_acc_ << "," << problem.dtype_c_ << "," - << problem.layout_a_ << "," << problem.layout_b_ << "," << problem.layout_c_ - << "," << problem.structured_sparsity_ << "," << name << "," << std::fixed - << std::setprecision(4) << perf.latency_ << "," << std::fixed - << std::setprecision(4) << perf.tflops_ << "," << std::fixed - << std::setprecision(4) << perf.bandwidth_ << "," << get_metric_name(metric) - << "\n"; - - if(!file) - { - std::cerr << "Warning: Error occurred while writing to CSV file." << std::endl; - } - } - } - - return kernel_instance; - } - - GemmProfiler(const GemmProfiler&) = delete; - GemmProfiler& operator=(const GemmProfiler&) = delete; - - private: - ~GemmProfiler() { kernel_instances_.clear(); } - GemmProfiler(Setting setting) : setting_(setting) {} - - Setting setting_; - - std::vector kernel_instances_; };