AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-378006 - Adding a new guarantee macro to support printing args.

Browse Source 
Change-Id: I2083ddaf8058759bdace7d3bb1e8e3670cbeddd2
This commit is contained in:
kjayapra-amd
2023-04-12 16:55:40 -04:00
committed by Karthik Jayaprakash
parent 82e4859e03
commit 7f41a96ffb
5 changed files with 32 additions and 37 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_code_object.cpp
+2 -4
View File
@@ -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());
hipamd/src/hip_code_object.hpp
+2 -3
View File
@@ -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:
hipamd/src/hip_fatbin.cpp
+13 -7
View File
@@ -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<hip::Devi
}
// At this line, image should be a valid ptr.
guarantee(image_ != nullptr, "Image cannot be nullptr, file did not map for some reason");
guarantee(image_ != nullptr, "Image cannot be nullptr, file:%s did not map for some reason",
fname_.c_str());
do {
@@ -121,7 +122,8 @@ hipError_t FatBinaryInfo::ExtractFatBinaryUsingCOMGR(const std::vector<hip::Devi
#if !defined(_WIN32)
// Using the file descriptor and file size, map the data object.
if (fdesc_ > 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::vector<hip::Devi
if ((comgr_status = amd_comgr_lookup_code_object(data_object, query_list_array,
unique_isa_names.size())) != AMD_COMGR_STATUS_SUCCESS) {
LogPrintfError("Setting data from file slice failed with status %d ", comgr_status);
hip_status = hipErrorInvalidValue;
hip_status = hipErrorInvalidValue;
break;
}
for (size_t isa_idx = 0; isa_idx < unique_isa_names.size(); ++isa_idx) {
auto unique_it = unique_isa_names.find(query_list_array[isa_idx].isa);
guarantee(unique_isa_names.cend() != unique_it, "Cannot find unique isa");
guarantee(unique_isa_names.cend() != unique_it, "Cannot find unique isa ");
unique_it->second = std::pair<size_t, size_t>
(static_cast<size_t>(query_list_array[isa_idx].size),
static_cast<size_t>(query_list_array[isa_idx].offset));
@@ -262,7 +264,11 @@ hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector<hip::Device*>& 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;
}
hipamd/src/hip_global.cpp
+4 -14
View File
@@ -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<cl_program>(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() {
hipamd/src/hip_platform.cpp
+11 -9
View File
@@ -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<hip::FatBinaryInfo**>(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,