mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-07-14 19:18:35 +00:00
523ec7c863
This change restructures the profiling process in Tile Engine into a base class for the Profiling and Problem structs. With this all files needed for Tile Engine will have a base struct and files in the gemm/ directory that can be extended for each GEMM variant. Only the Problem and Profiler structs along with the reference functions need to be defined. Profiling functions that are common to each operation have been moved into a common utility file.
255 lines
8.6 KiB
C++
255 lines
8.6 KiB
C++
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
|
|
// SPDX-License-Identifier: MIT
|
|
|
|
#pragma once
|
|
#include <iostream>
|
|
#include <functional>
|
|
#include <tuple>
|
|
#include <exception>
|
|
#include <sstream>
|
|
#include <vector>
|
|
#include <string>
|
|
|
|
#include "ck_tile/core.hpp"
|
|
#include "ck_tile/host.hpp"
|
|
|
|
// Helper function to determine if a layout is row-major
|
|
template <typename Layout>
|
|
constexpr auto is_row_major(Layout)
|
|
{
|
|
return ck_tile::bool_constant<std::is_same_v<Layout, ck_tile::tensor_layout::gemm::RowMajor>>{};
|
|
}
|
|
|
|
// 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 <typename Problem>
|
|
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 <typename ADataType, typename BDataType, typename AccDataType, typename CDataType>
|
|
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<sizeof(ADataType) < sizeof(BDataType), ADataType, BDataType>;
|
|
// Calculate thresholds
|
|
const auto rtol = ck_tile::get_relative_threshold<ComputeType, CDataType, AccDataType>(
|
|
ck_tile::integer_divide_ceil(K, kbatch));
|
|
const auto atol = ck_tile::get_absolute_threshold<ComputeType, CDataType, AccDataType>(
|
|
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<CDataType, CDataType, CDataType>(kbatch);
|
|
const auto atol_split_k = ck_tile::get_absolute_threshold<CDataType, CDataType, CDataType>(
|
|
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 <typename ADataType,
|
|
typename BDataType,
|
|
typename D0DataType,
|
|
typename AccDataType,
|
|
typename CDataType>
|
|
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<sizeof(ADataType) < sizeof(BDataType), ADataType, BDataType>;
|
|
|
|
using ComputeType =
|
|
std::conditional_t<sizeof(ComputeTypeAB) < sizeof(D0DataType), ComputeTypeAB, D0DataType>;
|
|
|
|
// Calculate thresholds
|
|
const auto rtol = ck_tile::get_relative_threshold<ComputeType, CDataType, AccDataType>(
|
|
ck_tile::integer_divide_ceil(K, kbatch));
|
|
|
|
const auto atol = ck_tile::get_absolute_threshold<ComputeType, CDataType, AccDataType>(
|
|
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<CDataType, CDataType, CDataType>(kbatch);
|
|
|
|
const auto atol_split_k = ck_tile::get_absolute_threshold<CDataType, CDataType, CDataType>(
|
|
max_accumulated_value, kbatch);
|
|
|
|
// Use higher threshold
|
|
return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k));
|
|
}
|