mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-07-17 09:08:35 +00:00
fix
This commit is contained in:
@@ -4,7 +4,7 @@
|
||||
execute_process(
|
||||
COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/gemm_instance_builder.py
|
||||
--working_path ${CMAKE_CURRENT_BINARY_DIR}
|
||||
# --config_json ${CMAKE_CURRENT_LIST_DIR}/configs/user_provided_config.json
|
||||
--config_json ${CMAKE_CURRENT_LIST_DIR}/configs/user_provided_config.json
|
||||
--list_blobs
|
||||
RESULT_VARIABLE ret
|
||||
)
|
||||
@@ -19,7 +19,7 @@ add_custom_command(
|
||||
OUTPUT ${GEMM_CODEGEN_BLOBS}
|
||||
COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_LIST_DIR}/gemm_instance_builder.py
|
||||
--working_path ${CMAKE_CURRENT_BINARY_DIR}
|
||||
# --config_json ${CMAKE_CURRENT_LIST_DIR}/configs/user_provided_config.json
|
||||
--config_json ${CMAKE_CURRENT_LIST_DIR}/configs/user_provided_config.json
|
||||
--gen_blobs
|
||||
)
|
||||
|
||||
|
||||
@@ -4,7 +4,7 @@ CK Tile Engine GEMM is used to generate and run GEMM kernels with different comb
|
||||
|
||||
# Kernel Configurations
|
||||
|
||||
ser can provide kernel configuration such as tile size, warp size, padding, pipeline, scheduler and epilogue in the config file. For reference please see `./configs/user_provided_config.json`. The Tile engine also has default kernel configuration to expand the range of kernel configuration which is saved in `./configs/default_config.json`.
|
||||
user can provide kernel configuration such as tile size, warp size, padding, pipeline, scheduler and epilogue in the config file. For reference please see `./configs/user_provided_config.json`. The Tile engine also has default kernel configuration to expand the range of kernel configuration which is saved in `./configs/default_config.json`.
|
||||
|
||||
## Build Instructions
|
||||
``` bash
|
||||
|
||||
@@ -7,8 +7,9 @@
|
||||
#include <vector>
|
||||
#include <filesystem>
|
||||
#include <memory>
|
||||
#include <fstream>
|
||||
#include <iomanip>
|
||||
|
||||
#include "ck/version.h"
|
||||
#include "ck/host_utility/device_prop.hpp"
|
||||
|
||||
struct GemmProblem
|
||||
@@ -57,7 +58,6 @@ struct PerformanceResult
|
||||
case Metric::BANDWIDTH: return a.bandwidth > b.bandwidth;
|
||||
default: throw std::invalid_argument("Unsupported metric type");
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
friend std::ostream& operator<<(std::ostream& os, const PerformanceResult& result)
|
||||
@@ -175,7 +175,8 @@ class GemmProfiler
|
||||
c_m_n_dev_result.SetZero();
|
||||
}
|
||||
|
||||
KernelInstance select_best_instance(Metric metric)
|
||||
KernelInstance select_best_instance(Metric metric,
|
||||
const std::string& csv_filename = "gemm_kernels.csv")
|
||||
{
|
||||
if(kernel_instances_.empty())
|
||||
throw std::runtime_error("Empty instances");
|
||||
@@ -192,6 +193,44 @@ class GemmProfiler
|
||||
<< "The best kernel instance is: " << kernel_instance << std::endl;
|
||||
std::cout << "**********************************" << std::endl;
|
||||
|
||||
if(!csv_filename.empty())
|
||||
{
|
||||
std::ofstream file(csv_filename, 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,"
|
||||
<< "latency(ms),tflops(TFlops),bandwidth(GB/s),metric\n";
|
||||
}
|
||||
|
||||
const auto& p = kernel_instance.problem;
|
||||
const auto& res = kernel_instance.perf_result;
|
||||
|
||||
file << get_rocm_version() << "," << ck::get_device_name() << "," << p.split_k
|
||||
<< "," << p.m << "," << p.n << "," << p.k << "," << p.stride_a << ","
|
||||
<< p.stride_b << "," << p.stride_c << "," << p.dtype_a << "," << p.dtype_b
|
||||
<< "," << p.dtype_acc << "," << p.dtype_c << "," << p.layout_a << ","
|
||||
<< p.layout_b << "," << p.layout_c << "," << std::fixed << std::setprecision(2)
|
||||
<< res.latency << "," << std::fixed << std::setprecision(2) << res.tflops
|
||||
<< "," << std::fixed << std::setprecision(2) << res.bandwidth << ","
|
||||
<< get_metric_name(metric) << "\n";
|
||||
|
||||
if(!file)
|
||||
{
|
||||
std::cerr << "Warning: Error occurred while writing to CSV file." << std::endl;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return kernel_instance;
|
||||
}
|
||||
|
||||
|
||||
@@ -35,7 +35,7 @@
|
||||
"tile_m": {
|
||||
"max": 512,
|
||||
"min": 256,
|
||||
"step": 256
|
||||
"step": 64
|
||||
},
|
||||
"tile_n": {
|
||||
"values": [
|
||||
|
||||
@@ -27,8 +27,8 @@ inline constexpr auto get_metric_name(Metric m)
|
||||
case Metric::LATENCY: return "latency";
|
||||
case Metric::TFLOPS: return "tflops";
|
||||
case Metric::BANDWIDTH: return "bandwidth";
|
||||
default: throw std::invalid_argument("Unsupported metric type");
|
||||
}
|
||||
return "unknown";
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
|
||||
Reference in New Issue
Block a user