From dc08a42fe05cdaea0c32f5f4ea49b22091b8ea6d Mon Sep 17 00:00:00 2001 From: Umang Yadav Date: Mon, 2 Oct 2023 18:17:58 +0000 Subject: [PATCH] just use one flush call --- .../gpu/device/include/migraphx/gpu/device/launch.hpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp b/src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp index fcc98c941a8..b36ef2b851f 100644 --- a/src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp +++ b/src/targets/gpu/device/include/migraphx/gpu/device/launch.hpp @@ -81,6 +81,14 @@ inline auto launch(hipStream_t stream, index_int global, index_int local) using f_type = decltype(f); dim3 nblocks(global / local); dim3 nthreads(local); + /* + hipGetLastError() returns error for the first failed HIP call that happened previously. + MIGraphX calls into various backend libraries and failed HIP calls can also happen there. + Calling hipGetLastError() would reset error code to hipSuccess, so that inside MIGraphX + failed call to hipLaunchKernelGGL() can be captured. + */ + hipError_t flush_call = hipGetLastError(); + (void)(flush_call); // cppcheck-suppress UseDeviceLaunch hipLaunchKernelGGL((launcher), nblocks, nthreads, 0, stream, f); hipError_t kernel_launch_status = hipGetLastError();