diff --git a/projects/clr/hipamd/src/hip_code_object.cpp b/projects/clr/hipamd/src/hip_code_object.cpp index cc26b415f5..800254b933 100644 --- a/projects/clr/hipamd/src/hip_code_object.cpp +++ b/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) { diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index fb630cfeed..c0ed09f004 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/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) { diff --git a/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp index 8512d684aa..4b16f7f605 100644 --- a/projects/clr/hipamd/src/hip_platform.cpp +++ b/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(const_cast(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(const_cast(hostFunction)); } constexpr auto gridDimYZmax = static_cast(std::numeric_limits::max()) + 1; diff --git a/projects/hip-tests/catch/hipTestMain/hip_test_context.cc b/projects/hip-tests/catch/hipTestMain/hip_test_context.cc index 1a321c7de0..6ea320f74b 100644 --- a/projects/hip-tests/catch/hipTestMain/hip_test_context.cc +++ b/projects/hip-tests/catch/hipTestMain/hip_test_context.cc @@ -123,11 +123,18 @@ std::string TestContext::getCurrentArch() { std::vector 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; })) {