diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index ae697775b5..5e169ee19f 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -625,10 +625,8 @@ hipError_t DynCO::initDynManagedVars(const std::string& managedVar) { } // Allocate managed memory for these symbols status = ihipMallocManaged(&pointer, dvar->size()); - if (status != hipSuccess) { - ClPrint(amd::LOG_ERROR, amd::LOG_API, "Status %d, failed to allocate managed memory", status); - guarantee(false, "Error during allocation of managed memory!"); - } + guarantee(status == hipSuccess, "Status %d, failed to allocate managed memory", status); + // update as manager variable and set managed memory pointer and size auto it = vars_.find(managedVar); it->second->setManagedVarInfo(pointer, dvar->size()); diff --git a/hipamd/src/hip_code_object.hpp b/hipamd/src/hip_code_object.hpp index db1225af39..1c784da934 100644 --- a/hipamd/src/hip_code_object.hpp +++ b/hipamd/src/hip_code_object.hpp @@ -104,9 +104,8 @@ public: } // Device ID Check to check if module is launched in the same device it was loaded. inline void CheckDeviceIdMatch() const { - if (device_id_ != ihipGetDevice()) { - guarantee(false, "Device mismatch from where this module is loaded"); - } + guarantee(device_id_ == ihipGetDevice(), "Device mismatch from where this module is loaded," + "device_id: %d ihipGetDevice:%d \n", device_id_, ihipGetDevice()); } private: diff --git a/hipamd/src/hip_fatbin.cpp b/hipamd/src/hip_fatbin.cpp index c2b7ff75b9..079b9a961b 100644 --- a/hipamd/src/hip_fatbin.cpp +++ b/hipamd/src/hip_fatbin.cpp @@ -36,10 +36,10 @@ FatBinaryInfo::~FatBinaryInfo() { if (fdesc_ > 0) { if (fsize_ && image_mapped_ && !amd::Os::MemoryUnmapFile(image_, fsize_)) { - guarantee(false, "Cannot unmap file"); + guarantee(false, "Cannot unmap file for fdesc: %d fsize: %d \n", fdesc_, fsize_); } if (!amd::Os::CloseFileHandle(fdesc_)) { - guarantee(false, "Cannot close file"); + guarantee(false, "Cannot close file for fdesc: %d \n", fdesc_); } } @@ -91,7 +91,8 @@ hipError_t FatBinaryInfo::ExtractFatBinaryUsingCOMGR(const std::vector 0) { - guarantee(fsize_ > 0, "Cannot have a file size of 0"); + guarantee(fsize_ > 0, "Cannot have a file size of 0, fdesc: %d fname: %s \n", + fdesc_, fname_.c_str()); if ((comgr_status = amd_comgr_set_data_from_file_slice(data_object, fdesc_, foffset_, fsize_)) != AMD_COMGR_STATUS_SUCCESS) { LogPrintfError("Setting data from file slice failed with status %d ", comgr_status); @@ -165,13 +167,13 @@ hipError_t FatBinaryInfo::ExtractFatBinaryUsingCOMGR(const std::vectorsecond = std::pair (static_cast(query_list_array[isa_idx].size), static_cast(query_list_array[isa_idx].offset)); @@ -262,7 +264,11 @@ hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector& devi } if (hip_error == hipErrorNoBinaryForGpu) { - LogPrintfError("hipErrorNoBinaryForGpu: Couldn't find binary for current devices! - %d",hip_error); + if (fname_.size() > 0) { + LogPrintfError("hipErrorNoBinaryForGpu: Couldn't find binary for file: %s", fname_.c_str()); + } else { + LogPrintfError("hipErrorNoBinaryForGpu: Couldn't find binary for ptr: 0x%x", image_); + } return hip_error; } diff --git a/hipamd/src/hip_global.cpp b/hipamd/src/hip_global.cpp index 5759cb715e..1042e1cdfa 100644 --- a/hipamd/src/hip_global.cpp +++ b/hipamd/src/hip_global.cpp @@ -50,14 +50,10 @@ DeviceVar::DeviceVar(std::string name, device::Program* dev_program = program->getDeviceProgram(*g_devices.at(deviceId)->devices()[0]); - if (dev_program == nullptr) { - LogPrintfError("Cannot get Device Program for module: 0x%x \n", hmod); - guarantee(false, "Cannot get Device Program"); - } + guarantee (dev_program != nullptr, "Cannot get Device Program for module: 0x%x \n", hmod); if(!dev_program->createGlobalVarObj(&amd_mem_obj_, &device_ptr_, &size_, name.c_str())) { - LogPrintfError("Cannot create Global Var obj for symbol: %s \n", name.c_str()); - guarantee(false, "Cannot create GlobalVar Obj"); + guarantee(false, "Cannot create GlobalVar Obj for symbol: %s \n", name.c_str()); } // Handle size 0 symbols @@ -93,16 +89,10 @@ DeviceFunc::DeviceFunc(std::string name, hipModule_t hmod) : dflock_("function l amd::Program* program = as_amd(reinterpret_cast(hmod)); const amd::Symbol *symbol = program->findSymbol(name.c_str()); - if (symbol == nullptr) { - LogPrintfError("Cannot find Symbol with name: %s \n", name.c_str()); - guarantee(false, "Cannot find Symbol"); - } + guarantee(symbol != nullptr, "Cannot find Symbol with name: %s \n", name.c_str()); kernel_ = new amd::Kernel(*program, *symbol, name); - if (kernel_ == nullptr) { - LogPrintfError("Cannot create kernel with name: %s \n", name.c_str()); - guarantee(false, "Cannot Create kernel"); - } + guarantee(kernel_ != nullptr, "Cannot Create kernel with name: %s \n", name.c_str()); } DeviceFunc::~DeviceFunc() { diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp index 6173f4d9c9..65a7551fca 100644 --- a/hipamd/src/hip_platform.cpp +++ b/hipamd/src/hip_platform.cpp @@ -88,7 +88,7 @@ extern "C" void __hipRegisterFunction(hip::FatBinaryInfo** modules, const void* hipError_t hip_error = hipSuccess; hip::Function* func = new hip::Function(std::string(deviceName), modules); hip_error = PlatformState::instance().registerStatFunction(hostFunction, func); - guarantee((hip_error == hipSuccess), "Cannot register Static function"); + guarantee((hip_error == hipSuccess), "Cannot register Static function, error: %d \n", hip_error); if (!enable_deferred_loading) { HIP_INIT_VOID(); @@ -96,7 +96,8 @@ extern "C" void __hipRegisterFunction(hip::FatBinaryInfo** modules, const void* for (size_t dev_idx = 0; dev_idx < g_devices.size(); ++dev_idx) { hip_error = PlatformState::instance().getStatFunc(&hfunc, hostFunction, dev_idx); - guarantee((hip_error == hipSuccess), "Cannot retrieve Static function"); + guarantee((hip_error == hipSuccess), "Cannot retrieve Static function, error: %d \n", + hip_error); } } } @@ -119,7 +120,7 @@ extern "C" void __hipRegisterVar( hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Variable, size, 0, 0, modules); hipError_t err = PlatformState::instance().registerStatGlobalVar(var, var_ptr); - guarantee((err == hipSuccess), "Cannot register Static Global Var"); + guarantee((err == hipSuccess), "Cannot register Static Global Var, error:%d \n", err); } extern "C" void __hipRegisterSurface( @@ -131,7 +132,7 @@ extern "C" void __hipRegisterSurface( hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Surface, sizeof(surfaceReference), 0, 0, modules); hipError_t err = PlatformState::instance().registerStatGlobalVar(var, var_ptr); - guarantee((err == hipSuccess), "Cannot register Static Glbal Var"); + guarantee((err == hipSuccess), "Cannot register Static Glbal Var, err:%d \n", err); } extern "C" void __hipRegisterManagedVar( @@ -147,17 +148,18 @@ extern "C" void __hipRegisterManagedVar( hip::Stream* stream = hip::getNullStream(); if (stream != nullptr) { status = ihipMemcpy(*pointer, init_value, size, hipMemcpyHostToDevice, *stream); - guarantee((status == hipSuccess), "Error during memcpy to managed memory!"); + guarantee((status == hipSuccess), "Error during memcpy to managed memory, error:%d \n!", + status); } else { ClPrint(amd::LOG_ERROR, amd::LOG_API, "Host Queue is NULL"); } } else { - guarantee(false, "Error during allocation of managed memory!"); + guarantee(false, "Error during allocation of managed memory!, error: %d \n", status); } hip::Var* var_ptr = new hip::Var(std::string(name), hip::Var::DeviceVarKind::DVK_Managed, pointer, size, align, reinterpret_cast(hipModule)); status = PlatformState::instance().registerStatManagedVar(var_ptr); - guarantee((status == hipSuccess), "Cannot register Static Managed Var"); + guarantee((status == hipSuccess), "Cannot register Static Managed Var, error: %d \n", status); } extern "C" void __hipRegisterTexture( @@ -169,12 +171,12 @@ extern "C" void __hipRegisterTexture( hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Texture, sizeof(textureReference), 0, 0, modules); hipError_t err = PlatformState::instance().registerStatGlobalVar(var, var_ptr); - guarantee((err == hipSuccess), "Cannot register Static Global Var"); + guarantee((err == hipSuccess), "Cannot register Static Global Var, status: %d \n", err); } extern "C" void __hipUnregisterFatBinary(hip::FatBinaryInfo** modules) { hipError_t err = PlatformState::instance().removeFatBinary(modules); - guarantee((err == hipSuccess), "Cannot Unregister Fat Binary"); + guarantee((err == hipSuccess), "Cannot Unregister Fat Binary, error:%d \n", err); } extern "C" hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem,