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

modified average finding to correct for repeats and changed 1000ms time limit to config parameter

Browse Source
This commit is contained in:
Harisankar Sadasivan
2024-04-17 16:46:28 -04:00
parent 4396a224d9
commit 5898ba83d9
2 changed files with 12 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
include/ck/host_utility/kernel_launch.hpp
View File

@@ -45,9 +45,10 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
float total_time = 0;
hip_check_error(hipEventElapsedTime(&total_time, start, stop));
total_time /= 10;
total_time /= stream_config.nrepeat_;
stream_config.cold_niters_ =
(1000.0 / total_time); // we need longer runtime to ramp up the clk on MI300s
(stream_config.time_limit_ms /
total_time); // we need longer runtime to ramp up the clk on MI300s
stream_config.nrepeat_ = stream_config.cold_niters_;
}
#endif
@@ -148,9 +149,10 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
float total_time = 0;
hip_check_error(hipEventElapsedTime(&total_time, start, stop));
total_time /= 10;
total_time /= stream_config.nrepeat_;
stream_config.cold_niters_ =
(1000.0 / total_time); // we need longer runtime to ramp up the clk on MI300s
(stream_config.nrepeat_ /
total_time); // we need longer runtime to ramp up the clk on MI300s
stream_config.nrepeat_ = stream_config.cold_niters_;
}
#endif

11
include/ck/stream_config.hpp
View File

@@ -8,9 +8,10 @@
struct StreamConfig
{
hipStream_t stream_id_ = nullptr;
bool time_kernel_ = false;
int log_level_ = 0;
mutable int cold_niters_ = 5;
mutable int nrepeat_ = 50;
hipStream_t stream_id_ = nullptr;
bool time_kernel_ = false;
int log_level_ = 0;
mutable int cold_niters_ = 5;
mutable int nrepeat_ = 50;
mutable int time_limit_ms = 1000; // for timing MI300 runs
};