ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-12 10:08:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add CK Logging

Browse Source
This commit is contained in:
Vsevolod Golovko
2025-08-26 17:40:12 +02:00
parent 84a7600bdc
commit fe98b0cd15
1 changed files with 19 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -38,7 +38,9 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
for(int i = 0; i < stream_config.cold_niters_; ++i)
{
kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
std::cout << "From CK 1:" << std::endl;
hip_check_error(hipGetLastError());
std::cout << "From CK 2:" << std::endl;
}
const int nrepeat = stream_config.nrepeat_;
@@ -48,39 +50,53 @@ float launch_and_time_kernel(const StreamConfig& stream_config,
}
hipEvent_t start, stop;
std::cout << "From CK 3:" << std::endl;
hip_check_error(hipEventCreate(&start));
std::cout << "From CK 4:" << std::endl;
hip_check_error(hipEventCreate(&stop));
std::cout << "From CK 5:" << std::endl;
hip_check_error(hipDeviceSynchronize());
std::cout << "From CK 6:" << std::endl;
hip_check_error(hipEventRecord(start, stream_config.stream_id_));
std::cout << "From CK 7:" << std::endl;
for(int i = 0; i < nrepeat; ++i)
{
kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
std::cout << "From CK 8:" << std::endl;
hip_check_error(hipGetLastError());
std::cout << "From CK 9:" << std::endl;
}
std::cout << "From CK 10:" << std::endl;
hip_check_error(hipEventRecord(stop, stream_config.stream_id_));
std::cout << "From CK 11:" << std::endl;
hip_check_error(hipEventSynchronize(stop));
std::cout << "From CK 12:" << std::endl;
float total_time = 0;
hip_check_error(hipEventElapsedTime(&total_time, start, stop));
std::cout << "From CK 13:" << std::endl;
hip_check_error(hipEventDestroy(start));
std::cout << "From CK 14:" << std::endl;
hip_check_error(hipEventDestroy(stop));
std::cout << "From CK 15:" << std::endl;
return total_time / nrepeat;
}
else
{
kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
std::cout << "From CK 16:" << std::endl;
hip_check_error(hipGetLastError());
return 0;
}
#else
kernel<<<grid_dim, block_dim, lds_byte, stream_config.stream_id_>>>(args...);
std::cout << "From CK 17:" << std::endl;
hip_check_error(hipGetLastError());
return 0;
@@ -96,6 +112,7 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
std::size_t lds_byte,
Args... args)
{
std::cout << "From CK 20:" << std::endl;
#if CK_TIME_KERNEL
if(stream_config.time_kernel_)
{
@@ -166,5 +183,6 @@ float launch_and_time_kernel_with_preprocess(const StreamConfig& stream_config,
return 0;
#endif
std::cout << "From CK 21:" << std::endl;
}
#endif