From 03d122e88f382ab8acb56d4f365f49f4bd808e75 Mon Sep 17 00:00:00 2001 From: Po Yen Chen Date: Wed, 26 Jun 2024 22:04:52 +0800 Subject: [PATCH] Replace hipDeviceSynchronize() by hipStreamSynchronize(stream) calls (#1359) [ROCm/composable_kernel commit: a32b1bc64794b8bd51f6924cb7d72dd1059803dd] --- include/ck_tile/host/timer.hpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/include/ck_tile/host/timer.hpp b/include/ck_tile/host/timer.hpp index e2baeaef7c..e5519643bf 100644 --- a/include/ck_tile/host/timer.hpp +++ b/include/ck_tile/host/timer.hpp @@ -27,7 +27,7 @@ struct gpu_timer CK_TILE_HOST void start(const hipStream_t& s) { - HIP_CHECK_ERROR(hipDeviceSynchronize()); + HIP_CHECK_ERROR(hipStreamSynchronize(s)); HIP_CHECK_ERROR(hipEventRecord(start_evt, s)); } @@ -51,15 +51,15 @@ struct gpu_timer struct cpu_timer { // torch.utils.benchmark.Timer(), there is a sync inside each timer callback - CK_TILE_HOST void start(const hipStream_t&) + CK_TILE_HOST void start(const hipStream_t& s) { - HIP_CHECK_ERROR(hipDeviceSynchronize()); + HIP_CHECK_ERROR(hipStreamSynchronize(s)); start_tick = std::chrono::high_resolution_clock::now(); } // torch.utils.benchmark.Timer(), there is a sync inside each timer callback - CK_TILE_HOST void stop(const hipStream_t&) + CK_TILE_HOST void stop(const hipStream_t& s) { - HIP_CHECK_ERROR(hipDeviceSynchronize()); + HIP_CHECK_ERROR(hipStreamSynchronize(s)); stop_tick = std::chrono::high_resolution_clock::now(); } // return in ms