From eea93d58a2b2a7086b3be5a7e977391a016afad6 Mon Sep 17 00:00:00 2001
From: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com>
Date: Tue, 9 Dec 2025 15:10:25 +0000
Subject: [PATCH] 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 | 8 +++++++-
 projects/clr/hipamd/src/hip_platform.cpp    | 6 ++++++
 2 files changed, 13 insertions(+), 1 deletion(-)

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