From a03682cb80219848385ba29f2dea9f86f9876f00 Mon Sep 17 00:00:00 2001 From: Thrupti Raj Lakshmana Gowda Date: Tue, 1 Jul 2025 12:59:49 -0500 Subject: [PATCH] Updating Runtime log for CK Tile Engine (#2431) * Updating runtime log message for CK TILE ENGINE * Fixing Clang Format * Update tile_engine/ops/gemm/README.md Co-authored-by: Aviral Goel --------- Co-authored-by: ThruptiRajLakshmanaGowda Co-authored-by: Aviral Goel --- tile_engine/ops/gemm/README.md | 3 ++- tile_engine/ops/gemm/benchmark_gemm.hpp | 8 +++++--- tile_engine/ops/gemm/gemm_profiler.hpp | 3 ++- 3 files changed, 9 insertions(+), 5 deletions(-) diff --git a/tile_engine/ops/gemm/README.md b/tile_engine/ops/gemm/README.md index 40cb9acd1c..e74da4b958 100644 --- a/tile_engine/ops/gemm/README.md +++ b/tile_engine/ops/gemm/README.md @@ -4,7 +4,8 @@ CK Tile Engine GEMM is used to generate and run GEMM kernels with different comb # Kernel Configurations -User can provide kernel configuration such as tile size, warp size, padding, pipeline, scheduler and epilogue in the config file with limited values. For reference please see `./configs/user_provided_config.json`. +Users can specify custom kernel configurations such as tile size, warp size, padding, pipeline, scheduler, and epilogue in the config file. This allows building only for selected configurations, significantly reducing build time. +For reference please see `./configs/user_provided_config.json`. The Tile engine also has a default kernel configuration for providing range of configuration parameter values, which helps users who lack kernel development experience to benchmark. For reference please see in `./configs/default_config.json` diff --git a/tile_engine/ops/gemm/benchmark_gemm.hpp b/tile_engine/ops/gemm/benchmark_gemm.hpp index 459a40b080..bbb9c1d715 100644 --- a/tile_engine/ops/gemm/benchmark_gemm.hpp +++ b/tile_engine/ops/gemm/benchmark_gemm.hpp @@ -163,7 +163,8 @@ auto calculate_rtol_atol(const ck_tile::index_t K, } /// @brief Function to compare the results of the device and host computations -bool compare(ck_tile::index_t K, +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) @@ -178,8 +179,9 @@ bool compare(ck_tile::index_t K, rtol_atol.at(ck_tile::number<0>{}), rtol_atol.at(ck_tile::number<1>{})); - std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{}) - << " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{}) << std::endl; + 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_profiler.hpp b/tile_engine/ops/gemm/gemm_profiler.hpp index 20f601d46e..2b0cbe7880 100644 --- a/tile_engine/ops/gemm/gemm_profiler.hpp +++ b/tile_engine/ops/gemm/gemm_profiler.hpp @@ -172,7 +172,8 @@ class GemmProfiler c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data()); bool verified_correct = !setting_.verify_ || - compare(gemm_problem.k_, gemm_problem.split_k_, c_m_n_dev_result, c_m_n_host_result); + compare( + name, gemm_problem.k_, gemm_problem.split_k_, c_m_n_dev_result, c_m_n_host_result); if(verified_correct) {