ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 00:58:44 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix threshold calculation

Browse Source
This commit is contained in:
Graner, Johannes
2026-01-28 09:18:03 -05:00
parent 55d8e9b4f0
commit 0eee2d3392
1 changed files with 6 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
profiler/include/profiler/profile_grouped_conv_bwd_weight_impl.hpp
View File

@@ -366,6 +366,11 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
const index_t num_accums = output.GetElementSize() / conv_param.K_;
const index_t num_accums_split_k = split_k_value;
// Get maximum accumulated value from reference
const std::size_t tensor_size =
weight_device_result.mDesc.GetElementSpaceSize();
max_accumulated_value =
gpu_reduce_max<WeiDataType>(gpu_ref_wei_buf.GetDeviceBuffer(), tensor_size);
// Calculate thresholds
auto rtol =
ck::utils::get_relative_threshold<ComputeType, WeiDataType, AccDataType>(
@@ -385,9 +390,7 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
rtol = std::max(rtol, rtol_split_k);
atol = std::max(atol, atol_split_k);
// Perform GPU verification (max value computed internally on GPU)
const std::size_t tensor_size =
weight_device_result.mDesc.GetElementSpaceSize();
// Perform GPU verification
auto gpu_result =
ck::profiler::gpu_verify<WeiDataType>(wei_device_buf.GetDeviceBuffer(),
gpu_ref_wei_buf.GetDeviceBuffer(),