diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index f2d7b38c6c..deeea225ad 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/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); diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index a3fa4919e7..83bd08be50 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/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(gridDim.x) * blockDim.x; size_t globalWorkSizeY = static_cast(gridDim.y) * blockDim.y; @@ -704,7 +707,7 @@ hipError_t hipLaunchCooperativeKernel_common(const void* f, dim3 gridDim, dim3 b globalWorkSizeY > std::numeric_limits::max() || globalWorkSizeZ > std::numeric_limits::max() || (blockDim.x * blockDim.y * blockDim.z > device->info().maxWorkGroupSize_)) { - return HIP_RETURN(hipErrorInvalidConfiguration); + return hipErrorInvalidConfiguration; } return ihipModuleLaunchKernel(func, static_cast(globalWorkSizeX), diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 6d085fe7f0..45ddd1665b 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/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