diff --git a/projects/clr/hipamd/src/hip_code_object.cpp b/projects/clr/hipamd/src/hip_code_object.cpp index 321f0975ea..cc26b415f5 100644 --- a/projects/clr/hipamd/src/hip_code_object.cpp +++ b/projects/clr/hipamd/src/hip_code_object.cpp @@ -450,12 +450,18 @@ hipError_t StatCO::getStatFunc(hipFunction_t* hfunc, const void* hostFunction, i // Lazy load FatBinaryInfo** module = it->second->moduleInfo(); - if (*(module) == nullptr) { + if (module != nullptr) { amd::ScopedLock lock(sclock_); if (*(module) == nullptr) { hipError_t err = digestFatBinary(module_to_hostModule_[module], *module); assert(err == hipSuccess); + if (err != hipSuccess) { + return err; + } } + } else { + // Module was nullptr + return hipErrorInvalidDeviceFunction; } return it->second->getStatFunc(hfunc, deviceId); diff --git a/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp index 3b7b56d0d2..8512d684aa 100644 --- a/projects/clr/hipamd/src/hip_platform.cpp +++ b/projects/clr/hipamd/src/hip_platform.cpp @@ -701,6 +701,12 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId); + + // Handle Invalid Image + if(hip_error == hipErrorInvalidImage) { + return hip_error; + } + if ((hip_error != hipSuccess) || (func == nullptr)) { // assume its hip function type if we did not get a valid output from static // func lookup