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

@@ -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 {                                   \