From 342c8cf01dab271b175d4bf62d3b650bc57f44cb Mon Sep 17 00:00:00 2001 From: "PoYen, Chen" Date: Mon, 24 Jun 2024 14:33:09 +0000 Subject: [PATCH] Call HIP_CHECK_ERROR() macro to get real source info --- include/ck_tile/host/kernel_launch.hpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/include/ck_tile/host/kernel_launch.hpp b/include/ck_tile/host/kernel_launch.hpp index e9c5a0c254..5c7bf12bfc 100644 --- a/include/ck_tile/host/kernel_launch.hpp +++ b/include/ck_tile/host/kernel_launch.hpp @@ -73,17 +73,17 @@ CK_TILE_HOST float launch_kernel(const stream_config& s, Callables... callables) { // clang-format off if(!s.time_kernel_) { - (callables(s),...); hip_check_error(hipGetLastError()); + (callables(s),...); HIP_CHECK_ERROR(hipGetLastError()); return 0; } if(s.is_gpu_timer_) { gpu_timer timer {}; // warmup - for(int i = 0; i < s.cold_niters_; i++) { (callables(s),...); } hip_check_error(hipGetLastError()); + for(int i = 0; i < s.cold_niters_; i++) { (callables(s),...); } HIP_CHECK_ERROR(hipGetLastError()); timer.start(s.stream_id_); - for(int i = 0; i < s.nrepeat_; i++) { (callables(s),...); } hip_check_error(hipGetLastError()); + for(int i = 0; i < s.nrepeat_; i++) { (callables(s),...); } HIP_CHECK_ERROR(hipGetLastError()); timer.stop(s.stream_id_); return timer.duration() / s.nrepeat_; @@ -92,10 +92,10 @@ CK_TILE_HOST float launch_kernel(const stream_config& s, Callables... callables) cpu_timer timer {}; // warmup - for(int i = 0; i < s.cold_niters_; i++) { (callables(s),...); } hip_check_error(hipGetLastError()); + for(int i = 0; i < s.cold_niters_; i++) { (callables(s),...); } HIP_CHECK_ERROR(hipGetLastError()); timer.start(s.stream_id_); - for(int i = 0; i < s.nrepeat_; i++) { (callables(s),...); } hip_check_error(hipGetLastError()); + for(int i = 0; i < s.nrepeat_; i++) { (callables(s),...); } HIP_CHECK_ERROR(hipGetLastError()); timer.stop(s.stream_id_); return timer.duration() / s.nrepeat_;