SWDEV-554626 - return correct error code (#1107)

* SWDEV-554626 - return hipErrorInvalidDeviceFunction when we can not load module
Return correct error code when modules are empty

* Match the error codes

* Revert the error code

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);

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