From 3dc47bc8392a6edb1d1a9a4b328b88c844f6f014 Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Thu, 9 Sep 2021 12:20:29 -0700 Subject: [PATCH] SWDEV-252801 - Fix hipGetLastError/hipPeekAtLastError hipGetLastError/hipPeekAtLastError should really return the last error produced by any of the runtime calls that did not successfuly complete, not just the error code of the last runtime call. For example, in this snippet of code: 1: hipLaunchKernelGGL(MyKernel, dim3(2), dim3(128), 0, 0, 10000,x); 2: hipDeviceSynchronize(); 3: hipError_t error = hipGetLastError(); If the call to hipLaunchKernelGGL fails, hipDeviceSynchronize (which succeeds) should not reset the last error to hipSuccess. hipGetLastError should still return the non-success error code returned by hipLaunchKernelGGL. The last error is reset to hipSuccess after calling hipGetLastError. Change-Id: Ib7c039067c53c94c99c8ecd83f54212bcef06f81 --- hipamd/src/hip_internal.hpp | 20 +++++++++++++------- 1 file changed, 13 insertions(+), 7 deletions(-) diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index 437a0c9d05..eccb87e0ce 100644 --- a/hipamd/src/hip_internal.hpp +++ b/hipamd/src/hip_internal.hpp @@ -97,16 +97,22 @@ typedef struct ihipIpcEventHandle_st { HIP_INIT() \ HIP_CB_SPAWNER_OBJECT(cid); -#define HIP_RETURN_DURATION(ret, ...) \ - hip::g_lastError = ret; \ +#define HIP_RETURN_DURATION(ret, ...) \ + hipError_t _lastError = (ret); \ + if (_lastError != hipSuccess) { \ + hip::g_lastError = _lastError; \ + } \ HIPPrintDuration(amd::LOG_INFO, amd::LOG_API, &startTimeUs, "%s: Returned %s : %s", \ - __func__, hipGetErrorName(hip::g_lastError), ToString( __VA_ARGS__ ).c_str()); \ - return hip::g_lastError; + __func__, hipGetErrorName(_lastError), ToString( __VA_ARGS__ ).c_str()); \ + return _lastError; #define HIP_RETURN(ret, ...) \ - hip::g_lastError = ret; \ - HIP_ERROR_PRINT(hip::g_lastError, __VA_ARGS__) \ - return hip::g_lastError; + hipError_t _lastError = (ret); \ + if (_lastError != hipSuccess) { \ + hip::g_lastError = _lastError; \ + } \ + HIP_ERROR_PRINT(_lastError, __VA_ARGS__) \ + return _lastError; #define HIP_RETURN_ONFAIL(func) \ do { \