From b36803746b7c72f1d2bf1666006d396b08e312ec 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 [ROCm/clr commit: 3dc47bc8392a6edb1d1a9a4b328b88c844f6f014] --- projects/clr/hipamd/src/hip_internal.hpp | 20 +++++++++++++------- 1 file changed, 13 insertions(+), 7 deletions(-) diff --git a/projects/clr/hipamd/src/hip_internal.hpp b/projects/clr/hipamd/src/hip_internal.hpp index 437a0c9d05..eccb87e0ce 100644 --- a/projects/clr/hipamd/src/hip_internal.hpp +++ b/projects/clr/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 { \