SWDEV-508225 remove assertions when loading fat binary (#2013)

* SWDEV-508225 - do not assert() after calling digestFatBinary() if it fails. Otherwise this causes assertions to trigger easily in systems that have an APU and a discrete GPU and the code was compiled for the discrete one

* SWDEV-508225 - fix that when using a non-existent ordinal in HIP_VISIBLE_DEVICES, getCurrentArch() would crash

projects/clr/hipamd/src/hip_code_object.cpp

@@ -454,7 +454,7 @@ hipError_t StatCO::getStatFunc(hipFunction_t* hfunc, const void* hostFunction, i
     amd::ScopedLock lock(sclock_);
     if (*(module) == nullptr) {
       hipError_t err = digestFatBinary(module_to_hostModule_[module], *module);
-      assert(err == hipSuccess);
+
       if (err != hipSuccess) {
         return err;
       }
@@ -479,8 +479,7 @@ hipError_t StatCO::getStatFuncAttr(hipFuncAttributes* func_attr, const void* hos
   // Lazy load
   FatBinaryInfo** module = it->second->moduleInfo();
   if (*(module) == nullptr) {
-    hipError_t err = digestFatBinary(module_to_hostModule_[module], *module);
-    assert(err == hipSuccess);
+    std::ignore = digestFatBinary(module_to_hostModule_[module], *module);
   }
 
   return it->second->getStatFuncAttr(func_attr, deviceId);
@@ -511,8 +510,7 @@ hipError_t StatCO::getStatGlobalVar(const void* hostVar, int deviceId, hipDevice
   // Lazy load
   FatBinaryInfo** module = it->second->moduleInfo();
   if (*(module) == nullptr) {
-    hipError_t err = digestFatBinary(module_to_hostModule_[module], *module);
-    assert(err == hipSuccess);
+    std::ignore = digestFatBinary(module_to_hostModule_[module], *module);
   }
 
   DeviceVar* dvar = nullptr;
@@ -538,8 +536,7 @@ hipError_t StatCO::initStatManagedVarDevicePtr(int deviceId) {
         // Lazy load
         FatBinaryInfo** module = var->moduleInfo();
         if (*(module) == nullptr) {
-          err = digestFatBinary(module_to_hostModule_[module], *module);
-          assert(err == hipSuccess);
+          std::ignore = digestFatBinary(module_to_hostModule_[module], *module);
         }
         hip::Stream* stream = g_devices.at(deviceId)->NullStream();
         if (stream == nullptr) {

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -1164,7 +1164,6 @@ class GraphKernelNode : public GraphNode {
     for (auto& command : commands_) {
       hipFunction_t func = getFunc(kernelParams_, dev_id_);
       hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func);
-      amd::Kernel* kernel = function->kernel();
       amd::ScopedLock lock(function->dflock_);
       command->enqueue();
       command->release();
@@ -1419,7 +1418,6 @@ class GraphKernelNode : public GraphNode {
       return hipErrorInvalidDeviceFunction;
     }
     hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func);
-    amd::Kernel* kernel = function->kernel();
     amd::ScopedLock lock(function->dflock_);
     status = validateKernelParams(&kernelParams_, func, dev_id_);
     if (hipSuccess != status) {

projects/clr/hipamd/src/hip_platform.cpp

@@ -365,9 +365,9 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
     int* maxBlocksPerCU, int* numBlocksPerGrid, int* bestBlockSize, const amd::Device& device,
     hipFunction_t func, int inputBlockSize, size_t dynamicSMemSize, bool bCalcPotentialBlkSz) {
   hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func);
-  const amd::Kernel& kernel = *function->kernel();
+  const amd::Kernel* kernel = function->kernel();
 
-  const device::Kernel::WorkGroupInfo* wrkGrpInfo = kernel.getDeviceKernel(device)->workGroupInfo();
+  const device::Kernel::WorkGroupInfo* wrkGrpInfo = kernel->getDeviceKernel(device)->workGroupInfo();
   if (bCalcPotentialBlkSz == false) {
     if (inputBlockSize <= 0) {
       return hipErrorInvalidValue;
@@ -702,15 +702,21 @@ 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) {
+  switch (hip_error) {
+  // invalid code object errors are propagated
+  case hipErrorInvalidKernelFile:
+  case hipErrorInvalidDeviceFunction:
+  case 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
-    func = reinterpret_cast<hipFunction_t>(const_cast<void *>(hostFunction));
+  case hipSuccess:
+    if (func) {
+      break;
+    }
+    // assume it is a hip function type if we did not get a valid output from static
+    // func lookup (i.e. if !func or hip_error != hipSuccess)
+    [[fallthrough]];
+  default:
+      func = reinterpret_cast<hipFunction_t>(const_cast<void *>(hostFunction));
   }
 
   constexpr auto gridDimYZmax = static_cast<uint64_t>(std::numeric_limits<uint16_t>::max()) + 1;

projects/hip-tests/catch/hipTestMain/hip_test_context.cc

@@ -123,11 +123,18 @@ std::string TestContext::getCurrentArch() {
     std::vector<std::string> filtered_archs;
     if (visible_devices.size() > 0) {
       for (size_t i = 0; i < visible_devices.size(); i++) {
-        filtered_archs.push_back(archs[visible_devices[i]]);
+        if (visible_devices[i] < archs.size()) {
+          filtered_archs.push_back(archs[visible_devices[i]]);
+        }
       }
     } else {
       filtered_archs = archs;
     }
+
+    if (filtered_archs.empty()) {
+      return "";
+    }
+
     auto first_filtered_arch = filtered_archs[0];
     if (!std::all_of(filtered_archs.begin(), filtered_archs.end(),
                      [&](const std::string& in) { return in == first_filtered_arch; })) {