ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 09:40:51 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add printouts to illustrate differences that can be seen on difference images

Browse Source
This commit is contained in:
Emily Martins
2025-10-20 15:33:01 +00:00
parent 86cf28d164
commit 50d984a353
2 changed files with 36 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
include/ck_tile/host/check_err.hpp
View File

@@ -117,7 +117,9 @@ CK_TILE_HOST double get_absolute_threshold(const double max_possible_num,
is_any_of<ComputeDataType, F8, BF8, F16, BF16, F32, pk_int4_t, I8, I32, int>::value,
"Warning: Unhandled ComputeDataType for setting up the absolute threshold!");
auto expo = std::log2(std::abs(max_possible_num));
auto expo = std::log2(std::abs(max_possible_num));
std::cout << "max_possible_num: " << max_possible_num << std::endl;
std::cout << "expo: " << expo << std::endl;
double compute_error = 0;
if constexpr(is_any_of<ComputeDataType, pk_int4_t, I8, I32, int>::value)
{

41
test/ck_tile/gemm/test_gemm_pipeline_util.hpp
View File

@@ -26,6 +26,7 @@ auto calculate_rtol_atol(const ck_tile::index_t K,
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
std::cout << "-------SPLIT K-------\n";
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>(
@@ -35,19 +36,24 @@ auto calculate_rtol_atol(const ck_tile::index_t K,
}
template <typename Type>
void print_tensor(int M, int N, ck_tile::HostTensor<Type>& tensor, std::string label)
void print_tensor(int M,
int N,
ck_tile::HostTensor<Type>& tensor,
std::string label,
bool use_type_convert = false)
{
int i = 0;
int i = 0;
int width = 7;
std::cout << "----------" << label << "----------\n";
for(int col = 0; col < N; ++col)
{
std::cout << std::setw(5) << col;
std::cout << std::setw(width) << col;
}
std::cout << std::endl;
for(int col = 0; col < N; ++col)
{
std::cout << std::setw(5) << "-----";
std::cout << std::setw(width) << "-----";
}
std::cout << std::endl;
@@ -55,8 +61,18 @@ void print_tensor(int M, int N, ck_tile::HostTensor<Type>& tensor, std::string l
{
for(int col = 0; col < N; ++col)
{
const double r = ck_tile::type_convert<float>(*std::next(std::begin(tensor), i));
std::cout << std::setw(5) << r;
float r;
if(use_type_convert)
{
r = ck_tile::type_convert<float>(*std::next(std::begin(tensor), i));
}
else
{
r = (*std::next(std::begin(tensor), i));
}
std::cout << std::setw(width) << r;
++i;
}
std::cout << std::endl;
@@ -443,14 +459,23 @@ class TestCkTileGemmPipeline : public ::testing::Test
ck_tile::reference_gemm<ADataType, BDataType, AccDataType, CDataType>(
a_m_k, b_k_n, c_m_n_host_ref);
print_tensor<CDataType>(M, N, c_m_n_host_ref, std::string{"REFERENCE"});
print_tensor<CDataType>(
M, N, c_m_n_host_ref, std::string{"REFERENCE: WITHOUT type_convert"});
std::cout << "\n\n\n";
print_tensor<CDataType>(M, N, c_m_n_dev_result, std::string{"GPU RESULT"});
print_tensor<CDataType>(
M, N, c_m_n_host_ref, std::string{"REFERENCE: WITH type_convert"}, true);
std::cout << "\n\n\n";
print_tensor<CDataType>(M, N, c_m_n_dev_result, std::string{"GPU RESULTS"});
const float max_accumulated_value =
*std::max_element(c_m_n_host_ref.mData.begin(), c_m_n_host_ref.mData.end());
std::cout << "max_accumulated_value: " << max_accumulated_value << std::endl;
const auto rtol_atol = calculate_rtol_atol<ADataType, BDataType, AccDataType, CDataType>(
K, kbatch, max_accumulated_value);
std::cout << "rtol: " << rtol_atol.at(ck_tile::number<0>{}) << std::endl;
std::cout << "atol: " << rtol_atol.at(ck_tile::number<1>{}) << std::endl;
pass = ck_tile::check_err(c_m_n_dev_result,
c_m_n_host_ref,
"Error: Incorrect results!",