mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-24 14:54:47 +00:00
9e3e865cec
* Fix splitK multiply_multiply_wp
* Add tests for gemm_multiply_multiply_wp
* Add tests for gemm_universal_preshuffle (KBatch = 1)
* Add tests gemm_blockscale_wp
* Fix splitk gemm universal preshuffle
* Run new tests on arch supporting fp8
* Restore example
* Fix strides profiler
* Fix tests
* Fix clang format
* Finalize profiler preshuffle with tolerances
* Minor improvements to splitk related changes
* Address review comments: clang format and ckProfiler typo
* Remove b_k_split_offset from SplitKBatchOffset struct
[ROCm/composable_kernel commit: 507d81c3af]
104 lines
2.4 KiB
C++
104 lines
2.4 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#pragma once
|
|
|
|
#include <type_traits>
|
|
#include "ck/utility/data_type.hpp"
|
|
|
|
namespace ck {
|
|
namespace profiler {
|
|
|
|
template <typename DataType, typename ComputeDataType = DataType>
|
|
inline __host__ __device__ constexpr double get_rtol()
|
|
{
|
|
if constexpr(std::is_same_v<DataType, float> && std::is_same_v<ComputeDataType, ck::tf32_t>)
|
|
{
|
|
return 1e-3;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, float>)
|
|
{
|
|
return 1e-3;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, double>)
|
|
{
|
|
return 1e-6;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::half_t>)
|
|
{
|
|
return 1e-3;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::bhalf_t>)
|
|
{
|
|
return 5e-2;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, int32_t>)
|
|
{
|
|
return 1e-1;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, int8_t>)
|
|
{
|
|
return 1e-1;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::f8_t>)
|
|
{
|
|
return 1e-1; // 240 and 224 are acceptable
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::bf8_t>)
|
|
{
|
|
return 1.5e-1; // 57344 and 49152 are acceptable
|
|
}
|
|
else
|
|
{
|
|
return 1e-3;
|
|
}
|
|
}
|
|
|
|
template <typename DataType, typename ComputeDataType = DataType>
|
|
inline __host__ __device__ constexpr double get_atol()
|
|
{
|
|
if constexpr(std::is_same_v<DataType, float> && std::is_same_v<ComputeDataType, ck::tf32_t>)
|
|
{
|
|
return 1e-3;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, float>)
|
|
{
|
|
return 1e-3;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, double>)
|
|
{
|
|
return 1e-6;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::half_t>)
|
|
{
|
|
return 1e-3;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::bhalf_t>)
|
|
{
|
|
return 5e-2;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, int32_t>)
|
|
{
|
|
return 1e-1;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, int8_t>)
|
|
{
|
|
return 1e-1;
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::f8_t>)
|
|
{
|
|
return 16.1; // 240 and 224 are acceptable
|
|
}
|
|
else if constexpr(std::is_same_v<DataType, ck::bf8_t>)
|
|
{
|
|
return 8192.1; // 57344 and 49152 are acceptable
|
|
}
|
|
else
|
|
{
|
|
return 1e-3;
|
|
}
|
|
}
|
|
|
|
} // namespace profiler
|
|
} // namespace ck
|