SWDEV-388278 - Fix return HIP_RETURN(x)

Change-Id: I940ee3660a2427eade8690a9318966fb04a1fbe8


[ROCm/clr commit: e6520ddf8a]

projects/clr/hipamd/src/hip_graph.cpp

@@ -2249,7 +2249,7 @@ hipError_t hipDeviceGetGraphMemAttribute(int device, hipGraphMemAttributeType at
     default:
       break;
   }
-  return HIP_RETURN(result);
+  HIP_RETURN(result);
 }
 
 // ================================================================================================
@@ -2274,7 +2274,7 @@ hipError_t hipDeviceSetGraphMemAttribute(int device, hipGraphMemAttributeType at
     default:
       break;
   }
-  return HIP_RETURN(result);
+  HIP_RETURN(result);
 }
 
 // ================================================================================================
@@ -2284,7 +2284,7 @@ hipError_t hipDeviceGraphMemTrim(int device) {
     HIP_RETURN(hipErrorInvalidDevice);
   }
   g_devices[device]->GetGraphMemoryPool()->TrimTo(0);
-  return HIP_RETURN(hipSuccess);
+  HIP_RETURN(hipSuccess);
 }
 
 // ================================================================================================
@@ -2417,7 +2417,7 @@ hipError_t hipGraphNodeSetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNod
   }
   if (!(hNode->GetType() == hipGraphNodeTypeKernel || hNode->GetType() == hipGraphNodeTypeMemcpy ||
         hNode->GetType() == hipGraphNodeTypeMemset)) {
-    return HIP_RETURN(hipErrorInvalidValue);
+    HIP_RETURN(hipErrorInvalidValue);
   }
   clonedNode->SetEnabled(isEnabled);
   HIP_RETURN(hipSuccess);
@@ -2436,7 +2436,7 @@ hipError_t hipGraphNodeGetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNod
   }
   if (!(hNode->GetType() == hipGraphNodeTypeKernel || hNode->GetType() == hipGraphNodeTypeMemcpy ||
         hNode->GetType() == hipGraphNodeTypeMemset)) {
-    return HIP_RETURN(hipErrorInvalidValue);
+    HIP_RETURN(hipErrorInvalidValue);
   }
   *isEnabled = clonedNode->GetEnabled();
   HIP_RETURN(hipSuccess);

projects/clr/hipamd/src/hip_module.cpp

@@ -690,12 +690,15 @@ hipError_t hipLaunchCooperativeKernel_common(const void* f, dim3 gridDim, dim3 b
                                              void** kernelParams, uint32_t sharedMemBytes,
                                              hipStream_t hStream) {
   if (!hip::isValid(hStream)) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
 
   hipFunction_t func = nullptr;
   int deviceId = hip::Stream::DeviceId(hStream);
-  HIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, f, deviceId));
+  hipError_t getStatFuncError = PlatformState::instance().getStatFunc(&func, f, deviceId);
+  if (getStatFuncError != hipSuccess) {
+    return getStatFuncError;
+  }
   const amd::Device* device = g_devices[deviceId]->devices()[0];
   size_t globalWorkSizeX = static_cast<size_t>(gridDim.x) * blockDim.x;
   size_t globalWorkSizeY = static_cast<size_t>(gridDim.y) * blockDim.y;
@@ -704,7 +707,7 @@ hipError_t hipLaunchCooperativeKernel_common(const void* f, dim3 gridDim, dim3 b
       globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeZ > std::numeric_limits<uint32_t>::max() ||
       (blockDim.x * blockDim.y * blockDim.z > device->info().maxWorkGroupSize_)) {
-    return HIP_RETURN(hipErrorInvalidConfiguration);
+    return hipErrorInvalidConfiguration;
   }
 
   return ihipModuleLaunchKernel(func, static_cast<uint32_t>(globalWorkSizeX),

projects/clr/hipamd/src/hip_stream.cpp

@@ -805,7 +805,7 @@ hipError_t hipStreamGetDevice(hipStream_t stream, hipDevice_t* device) {
   }
 
   if (!hip::isValid(stream)) {
-    return HIP_RETURN(hipErrorContextIsDestroyed);
+    HIP_RETURN(hipErrorContextIsDestroyed);
   }
 
   if (stream == nullptr) {  // handle null stream