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: 3dc47bc839]

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