SWDEV-368881 - Revert enable hipLaunchCooperativeKernel for hipRTC compiled function

This reverts commit f870f04dd9.

Reason for revert: Doesn't match with CUDA behavior

Change-Id: I413f3d241f864c1e7c21681c5fff8f216ef54306


[ROCm/clr commit: 4f5d180ca7]

projects/clr/hipamd/src/hip_code_object.cpp

@@ -562,20 +562,6 @@ hipError_t DynCO::getDynFunc(hipFunction_t* hfunc, std::string func_name) {
   return it->second->getDynFunc(hfunc, module());
 }
 
-bool DynCO::isValidDynFunc(const hipFunction_t& hfunc) {
-  amd::ScopedLock lock(dclock_);
-
-  CheckDeviceIdMatch();
-
-  for (auto it : functions_) {
-    if (it.second->isValidDynFunc(hfunc)) {
-      return true;
-    }
-  }
-
-  return false;
-}
-
 hipError_t DynCO::initDynManagedVars(const std::string& managedVar) {
   amd::ScopedLock lock(dclock_);
   DeviceVar* dvar;

projects/clr/hipamd/src/hip_code_object.hpp

@@ -92,9 +92,6 @@ public:
   hipError_t getDynFunc(hipFunction_t* hfunc, std::string func_name);
   hipError_t getDeviceVar(DeviceVar** dvar, std::string var_name);
 
-  // Verify whether the dynamic function is present or not
-  bool isValidDynFunc(const hipFunction_t& hfunc);
-
   hipError_t getManagedVarPointer(std::string name, void** pointer, size_t* size_ptr) const {
     auto it = vars_.find(name);
     if (it != vars_.end() && it->second->getVarKind() == Var::DVK_Managed) {

projects/clr/hipamd/src/hip_global.cpp

@@ -135,14 +135,6 @@ hipError_t Function::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod) {
   return hipSuccess;
 }
 
-bool Function::isValidDynFunc(const hipFunction_t& hfunc) {
-  guarantee((dFunc_.size() == g_devices.size()), "dFunc Size mismatch");
-  if (dFunc_[ihipGetDevice()] == nullptr) {
-    return false;
-  }
-  return (dFunc_[ihipGetDevice()]->asHipFunction() == hfunc);
-}
-
 hipError_t Function::getStatFunc(hipFunction_t* hfunc, int deviceId) {
   guarantee(modules_ != nullptr, "Module not initialized");
 

projects/clr/hipamd/src/hip_global.hpp

@@ -62,8 +62,6 @@ public:
 
   //Return DeviceFunc for this this dynamically loaded module
   hipError_t getDynFunc(hipFunction_t* hfunc, hipModule_t hmod);
-  // Verify whether the dynamic device function is present or not
-  bool isValidDynFunc(const hipFunction_t& hfunc);
 
   //Return Device Func & attr . Generate/build if not already done so.
   hipError_t getStatFunc(hipFunction_t *hfunc, int deviceId);

projects/clr/hipamd/src/hip_module.cpp

@@ -497,15 +497,7 @@ hipError_t hipLaunchCooperativeKernel_common(const void* f, dim3 gridDim, dim3 b
 
   hipFunction_t func = nullptr;
   int deviceId = hip::Stream::DeviceId(hStream);
-  hipError_t status = PlatformState::instance().getStatFunc(&func, f, deviceId);
-  if (status != hipSuccess) {
-    // Check if its a dynamic function
-    if (!PlatformState::instance().isValidDynFunc(hipFunction_t(f))) {
-      return status;
-    }
-    func = (hipFunction_t)f;
-  }
-
+  HIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, f, deviceId));
   size_t globalWorkSizeX = static_cast<size_t>(gridDim.x) * blockDim.x;
   size_t globalWorkSizeY = static_cast<size_t>(gridDim.y) * blockDim.y;
   size_t globalWorkSizeZ = static_cast<size_t>(gridDim.z) * blockDim.z;

projects/clr/hipamd/src/hip_platform.cpp

@@ -747,18 +747,6 @@ hipError_t PlatformState::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod,
   return it->second->getDynFunc(hfunc, func_name);
 }
 
-bool PlatformState::isValidDynFunc(const hipFunction_t& hfunc) {
-  amd::ScopedLock lock(lock_);
-
-  for (auto it : dynCO_map_) {
-    if (it.second->isValidDynFunc(hfunc)) {
-      return true;
-    }
-  }
-
-  return false;
-}
-
 hipError_t PlatformState::getDynGlobalVar(const char* hostVar, hipModule_t hmod,
                                           hipDeviceptr_t* dev_ptr, size_t* size_ptr) {
   amd::ScopedLock lock(lock_);

projects/clr/hipamd/src/hip_platform.hpp

@@ -47,8 +47,6 @@ class PlatformState {
   hipError_t unloadModule(hipModule_t hmod);
 
   hipError_t getDynFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* func_name);
-  // Verify whether the dynamic function is present or not
-  bool isValidDynFunc(const hipFunction_t& hfunc);
   hipError_t getDynGlobalVar(const char* hostVar, hipModule_t hmod, hipDeviceptr_t* dev_ptr,
                              size_t* size_ptr);
   hipError_t getDynTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef);