From 086aee0fcb32c1428809df8c781b1b64c1296382 Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Wed, 2 Mar 2022 11:46:56 -0800 Subject: [PATCH] SWDEV-301667 - Fix build warnings Change-Id: Ice23b3d1a19a7958ca5c3caff34db2934b361a98 --- hipamd/src/hip_code_object.cpp | 27 ++++++++++------- hipamd/src/hip_device_runtime.cpp | 24 +++++++-------- hipamd/src/hip_global.cpp | 2 +- hipamd/src/hip_module.cpp | 13 ++++---- hipamd/src/hip_platform.cpp | 49 +++++++++++++++++++------------ hipamd/src/hip_stream.cpp | 6 ++-- hipamd/src/hip_surface.cpp | 5 ++-- hipamd/src/hip_texture.cpp | 4 +-- 8 files changed, 76 insertions(+), 54 deletions(-) diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index 843ea24591..4595e92ba0 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -518,7 +518,8 @@ DynCO::~DynCO() { for (auto& elem : vars_) { if(elem.second->getVarKind() == Var::DVK_Managed) { - ihipFree(elem.second->getManagedVarPtr()); + hipError_t err = ihipFree(elem.second->getManagedVarPtr()); + assert(err == hipSuccess); } delete elem.second; } @@ -543,8 +544,8 @@ hipError_t DynCO::getDeviceVar(DeviceVar** dvar, std::string var_name) { return hipErrorNotFound; } - it->second->getDeviceVar(dvar, device_id_, module()); - return hipSuccess; + hipError_t err = it->second->getDeviceVar(dvar, device_id_, module()); + return err; } hipError_t DynCO::getDynFunc(hipFunction_t* hfunc, std::string func_name) { @@ -623,6 +624,7 @@ hipError_t DynCO::initDynManagedVars(const std::string& managedVar) { hipError_t DynCO::populateDynGlobalVars() { amd::ScopedLock lock(dclock_); + hipError_t err = hipSuccess; std::vector var_names; std::string managedVarExt = ".managed"; // For Dynamic Modules there is only one hipFatBinaryDevInfo_ @@ -643,10 +645,10 @@ hipError_t DynCO::populateDynGlobalVars() { if (elem.find(managedVarExt) != std::string::npos) { std::string managedVar = elem; managedVar.erase(managedVar.length() - managedVarExt.length(), managedVarExt.length()); - initDynManagedVars(managedVar); + err = initDynManagedVars(managedVar); } } - return hipSuccess; + return err; } hipError_t DynCO::populateDynGlobalFuncs() { @@ -706,7 +708,8 @@ FatBinaryInfo** StatCO::addFatBinary(const void* data, bool initialized) { amd::ScopedLock lock(sclock_); if (initialized) { - digestFatBinary(data, modules_[data]); + hipError_t err = digestFatBinary(data, modules_[data]); + assert(err == hipSuccess); } return &modules_[data]; } @@ -727,7 +730,8 @@ hipError_t StatCO::removeFatBinary(FatBinaryInfo** module) { auto it = managedVars_.begin(); while (it != managedVars_.end()) { if ((*it)->moduleInfo() == module) { - ihipFree((*it)->getManagedVarPtr()); + hipError_t err = ihipFree((*it)->getManagedVarPtr()); + assert(err == hipSuccess); delete *it; it = managedVars_.erase(it); } else { @@ -826,7 +830,7 @@ hipError_t StatCO::registerStatManagedVar(Var* var) { hipError_t StatCO::initStatManagedVarDevicePtr(int deviceId) { amd::ScopedLock lock(sclock_); - + hipError_t err = hipSuccess; if (managedVarsDevicePtrInitalized_.find(deviceId) == managedVarsDevicePtrInitalized_.end() || !managedVarsDevicePtrInitalized_[deviceId]) { for (auto var : managedVars_) { @@ -835,8 +839,9 @@ hipError_t StatCO::initStatManagedVarDevicePtr(int deviceId) { amd::HostQueue* queue = hip::getNullStream(); if(queue != nullptr) { - ihipMemcpy(reinterpret_cast
(dvar->device_ptr()), var->getManagedVarPtr(), - dvar->size(), hipMemcpyHostToDevice, *queue); + err = ihipMemcpy(reinterpret_cast
(dvar->device_ptr()), + var->getManagedVarPtr(), + dvar->size(), hipMemcpyHostToDevice, *queue); } else { ClPrint(amd::LOG_ERROR, amd::LOG_API, "Host Queue is NULL"); return hipErrorInvalidResourceHandle; @@ -844,6 +849,6 @@ hipError_t StatCO::initStatManagedVarDevicePtr(int deviceId) { } managedVarsDevicePtrInitalized_[deviceId] = true; } - return hipSuccess; + return err; } }; //namespace: hip diff --git a/hipamd/src/hip_device_runtime.cpp b/hipamd/src/hip_device_runtime.cpp index 0e9733d94e..642e40092f 100644 --- a/hipamd/src/hip_device_runtime.cpp +++ b/hipamd/src/hip_device_runtime.cpp @@ -33,7 +33,7 @@ hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) { *device = 0; cl_uint maxMatchedCount = 0; int count = 0; - ihipDeviceGetCount(&count); + HIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); for (cl_int i = 0; i< count; ++i) { hipDeviceProp_t currentProp = {0}; @@ -145,17 +145,15 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) } int count = 0; - ihipDeviceGetCount(&count); + HIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); + if (device < 0 || device >= count) { HIP_RETURN(hipErrorInvalidDevice); } //FIXME: should we cache the props, or just select from deviceHandle->info_? hipDeviceProp_t prop = {0}; - hipError_t err = ihipGetDeviceProperties(&prop, device); - if (err != hipSuccess) { - HIP_RETURN(err); - } + HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, device)); switch (attr) { case hipDeviceAttributeMaxThreadsPerBlock: @@ -344,12 +342,12 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) { reinterpret_cast(&pciBusID), reinterpret_cast(&pciDeviceID)) == 0x3) { int count = 0; - ihipDeviceGetCount(&count); + HIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); for (cl_int i = 0; i < count; i++) { hipDevice_t dev; - ihipDeviceGet(&dev, i); hipDeviceProp_t prop; - ihipGetDeviceProperties(&prop, dev); + HIP_RETURN_ONFAIL(ihipDeviceGet(&dev, i)); + HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, dev)); if ((pciBusID == prop.pciBusID) && (pciDomainID == prop.pciDomainID) && (pciDeviceID == prop.pciDeviceID)) { @@ -387,7 +385,7 @@ hipError_t hipDeviceGetLimit ( size_t* pValue, hipLimit_t limit ) { } if(limit == hipLimitMallocHeapSize) { hipDeviceProp_t prop; - ihipGetDeviceProperties(&prop, ihipGetDevice()); + HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, ihipGetDevice())); *pValue = prop.totalGlobalMem; HIP_RETURN(hipSuccess); @@ -401,7 +399,8 @@ hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) { HIP_INIT_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device); int count; - ihipDeviceGetCount(&count); + HIP_RETURN_ONFAIL(ihipDeviceGetCount(&count)); + if (device < 0 || device >= count) { HIP_RETURN(hipErrorInvalidDevice); } @@ -411,8 +410,7 @@ hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) { } hipDeviceProp_t prop; - ihipGetDeviceProperties(&prop, device); - + HIP_RETURN_ONFAIL(ihipGetDeviceProperties(&prop, device)); snprintf (pciBusId, len, "%04x:%02x:%02x.0", prop.pciDomainID, prop.pciBusID, diff --git a/hipamd/src/hip_global.cpp b/hipamd/src/hip_global.cpp index 0d9c2b44ab..eb56b6360b 100644 --- a/hipamd/src/hip_global.cpp +++ b/hipamd/src/hip_global.cpp @@ -62,7 +62,7 @@ DeviceVar::~DeviceVar() { if (shadowVptr != nullptr) { textureReference* texRef = reinterpret_cast(shadowVptr); - ihipUnbindTexture(texRef); + hipError_t err = ihipUnbindTexture(texRef); delete texRef; shadowVptr = nullptr; } diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index a4f1fbfc70..744283b024 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -265,9 +265,12 @@ hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX, int max_blocks_per_grid = 0; int best_block_size = 0; int block_size = blockDimX * blockDimY * blockDimZ; - hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks, &max_blocks_per_grid, + hipError_t err = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks, &max_blocks_per_grid, &best_block_size, *device, f, block_size, sharedMemBytes, true); + if (err != hipSuccess) { + return err; + } if (((globalWorkSizeX * globalWorkSizeY * globalWorkSizeZ) / block_size) > unsigned(max_blocks_per_grid)) { return hipErrorCooperativeLaunchTooLarge; @@ -544,13 +547,13 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL int numDevices, unsigned int flags, uint32_t extFlags) { int numActiveGPUs = 0; - ihipDeviceGetCount(&numActiveGPUs); + hipError_t result = hipSuccess; + result = ihipDeviceGetCount(&numActiveGPUs); if ((numDevices > numActiveGPUs) || (launchParamsList == nullptr)) { return hipErrorInvalidValue; } - hipError_t result = hipErrorUnknown; uint64_t allGridSize = 0; std::vector mgpu_list(numDevices); @@ -683,7 +686,7 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const // have the default read mode set to normalized float. (*texRef)->readMode = hipReadModeNormalizedFloat; - PlatformState::instance().registerTexRef(*texRef, hmod, std::string(name)); + hipError_t err = PlatformState::instance().registerTexRef(*texRef, hmod, std::string(name)); - HIP_RETURN(hipSuccess); + HIP_RETURN(err); } diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp index bcae85601c..26de428030 100644 --- a/hipamd/src/hip_platform.cpp +++ b/hipamd/src/hip_platform.cpp @@ -99,17 +99,18 @@ extern "C" void __hipRegisterFunction( char *var = getenv("HIP_ENABLE_DEFERRED_LOADING"); return var ? atoi(var) : 1; }() }; - + hipError_t hip_error = hipSuccess; hip::Function* func = new hip::Function(std::string(deviceName), modules); - PlatformState::instance().registerStatFunction(hostFunction, func); + hip_error = PlatformState::instance().registerStatFunction(hostFunction, func); + guarantee((hip_error == hipSuccess), "Cannot register Static function"); if (!enable_deferred_loading) { HIP_INIT_VOID(); hipFunction_t hfunc = nullptr; - hipError_t hip_error = hipSuccess; + 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"); } } } @@ -130,7 +131,8 @@ extern "C" void __hipRegisterVar( int global) // Unknown, always 0 { hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Variable, size, 0, 0, modules); - PlatformState::instance().registerStatGlobalVar(var, var_ptr); + hipError_t err = PlatformState::instance().registerStatGlobalVar(var, var_ptr); + guarantee((err == hipSuccess), "Cannot register Static Global Var"); } extern "C" void __hipRegisterSurface(hip::FatBinaryInfo** modules, // The device modules containing code object @@ -139,7 +141,8 @@ extern "C" void __hipRegisterSurface(hip::FatBinaryInfo** modules, // The d char* deviceVar, // Variable name in device code int type, int ext) { hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Surface, sizeof(surfaceReference), 0, 0, modules); - PlatformState::instance().registerStatGlobalVar(var, var_ptr); + hipError_t err = PlatformState::instance().registerStatGlobalVar(var, var_ptr); + guarantee((err == hipSuccess), "Cannot register Static Glbal Var"); } extern "C" void __hipRegisterManagedVar(void *hipModule, // Pointer to hip module returned from __hipRegisterFatbinary @@ -154,7 +157,8 @@ extern "C" void __hipRegisterManagedVar(void *hipModule, // Pointer to hip mod if( status == hipSuccess) { amd::HostQueue* queue = hip::getNullStream(); if(queue != nullptr) { - ihipMemcpy(*pointer, init_value, size, hipMemcpyHostToDevice, *queue); + status = ihipMemcpy(*pointer, init_value, size, hipMemcpyHostToDevice, *queue); + guarantee((status == hipSuccess), "Error during memcpy to managed memory!"); } else { ClPrint(amd::LOG_ERROR, amd::LOG_API, "Host Queue is NULL"); } @@ -163,7 +167,8 @@ extern "C" void __hipRegisterManagedVar(void *hipModule, // Pointer to hip mod } hip::Var* var_ptr = new hip::Var(std::string(name), hip::Var::DeviceVarKind::DVK_Managed, pointer, size, align, reinterpret_cast(hipModule)); - PlatformState::instance().registerStatManagedVar(var_ptr); + status = PlatformState::instance().registerStatManagedVar(var_ptr); + guarantee((status == hipSuccess), "Cannot register Static Managed Var"); } extern "C" void __hipRegisterTexture(hip::FatBinaryInfo** modules, // The device modules containing code object @@ -172,12 +177,14 @@ extern "C" void __hipRegisterTexture(hip::FatBinaryInfo** modules, // The d char* deviceVar, // Variable name in device code int type, int norm, int ext) { hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Texture, sizeof(textureReference), 0, 0, modules); - PlatformState::instance().registerStatGlobalVar(var, var_ptr); + hipError_t err = PlatformState::instance().registerStatGlobalVar(var, var_ptr); + guarantee((err == hipSuccess), "Cannot register Static Global Var"); } extern "C" void __hipUnregisterFatBinary(hip::FatBinaryInfo** modules) { - PlatformState::instance().removeFatBinary(modules); + hipError_t err = PlatformState::instance().removeFatBinary(modules); + guarantee((err == hipSuccess), "Cannot Unregister Fat Binary"); } extern "C" hipError_t hipConfigureCall( @@ -603,15 +610,19 @@ void hipLaunchKernelGGLImpl( } hipFunction_t func = nullptr; - hipError_t hip_error = PlatformState::instance().getStatFunc(&func, reinterpret_cast(function_address), deviceId); + hipError_t hip_error = + PlatformState::instance().getStatFunc(&func, + reinterpret_cast(function_address), + deviceId); if ((hip_error != hipSuccess) || (func == nullptr)) { LogPrintfError("Cannot find the static function: 0x%x", function_address); } - hipModuleLaunchKernel(func, - numBlocks.x, numBlocks.y, numBlocks.z, - dimBlocks.x, dimBlocks.y, dimBlocks.z, - sharedMemBytes, stream, nullptr, kernarg); + hip_error = hipModuleLaunchKernel(func, + numBlocks.x, numBlocks.y, numBlocks.z, + dimBlocks.x, dimBlocks.y, dimBlocks.z, + sharedMemBytes, stream, nullptr, kernarg); + assert(hip_error == hipSuccess); } void hipLaunchCooperativeKernelGGLImpl( @@ -624,8 +635,9 @@ void hipLaunchCooperativeKernelGGLImpl( { HIP_INIT_VOID(); - hipLaunchCooperativeKernel(reinterpret_cast(function_address), + hipError_t err = hipLaunchCooperativeKernel(reinterpret_cast(function_address), numBlocks, dimBlocks, kernarg, sharedMemBytes, stream); + assert(err == hipSuccess); } } @@ -730,7 +742,8 @@ void PlatformState::init() } initialized_ = true; for (auto& it : statCO_.modules_) { - digestFatBinary(it.first, it.second); + hipError_t err = digestFatBinary(it.first, it.second); + assert(err == hipSuccess); } for (auto &it : statCO_.vars_) { it.second->resize_dVar(g_devices.size()); @@ -819,7 +832,7 @@ hipError_t PlatformState::getDynGlobalVar(const char* hostVar, hipModule_t hmod, return hipErrorNotFound; } *dev_ptr = nullptr; - it->second->getManagedVarPointer(hostVar, dev_ptr, size_ptr); + IHIP_RETURN_ONFAIL(it->second->getManagedVarPointer(hostVar, dev_ptr, size_ptr)); // if dev_ptr is nullptr, hostvar is not in managed variable list if (*dev_ptr == nullptr) { hip::DeviceVar* dvar = nullptr; diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index d3720526b8..1fecfe45ab 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -62,7 +62,8 @@ hipError_t Stream::EndCapture() { } for (auto stream : parallelCaptureStreams_) { hip::Stream* s = reinterpret_cast(stream); - s->EndCapture(); + hipError_t err = s->EndCapture(); + assert(err == hipSuccess); } captureStatus_ = hipStreamCaptureStatusNone; pCaptureGraph_ = nullptr; @@ -307,7 +308,8 @@ public: // There is a scenario where hipResetDevice destroys stream per thread // hence isValid check is required to make sure only valid stream is used if (m_stream == nullptr || !hip::isValid(m_stream)) { - ihipStreamCreate(&m_stream, hipStreamDefault, hip::Stream::Priority::Normal); + hipError_t err = ihipStreamCreate(&m_stream, hipStreamDefault, hip::Stream::Priority::Normal); + assert(err == hipSuccess); } return m_stream; } diff --git a/hipamd/src/hip_surface.cpp b/hipamd/src/hip_surface.cpp index a8e7f71d2c..3ec495b668 100644 --- a/hipamd/src/hip_surface.cpp +++ b/hipamd/src/hip_surface.cpp @@ -64,8 +64,9 @@ hipError_t ihipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, image = as_amd(memObj)->asImage(); void* surfObjectBuffer = nullptr; - ihipMalloc(&surfObjectBuffer, sizeof(__hip_surface), CL_MEM_SVM_FINE_GRAIN_BUFFER); - if (surfObjectBuffer == nullptr) { + hipError_t err = ihipMalloc(&surfObjectBuffer, sizeof(__hip_surface), + CL_MEM_SVM_FINE_GRAIN_BUFFER); + if (surfObjectBuffer == nullptr || err != hipSuccess) { return hipErrorOutOfMemory; } *pSurfObject = new (surfObjectBuffer) __hip_surface{image, *pResDesc}; diff --git a/hipamd/src/hip_texture.cpp b/hipamd/src/hip_texture.cpp index bc73f731ed..34e9787300 100644 --- a/hipamd/src/hip_texture.cpp +++ b/hipamd/src/hip_texture.cpp @@ -297,8 +297,8 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, } void *texObjectBuffer = nullptr; - ihipMalloc(&texObjectBuffer, sizeof(__hip_texture), CL_MEM_SVM_FINE_GRAIN_BUFFER); - if (texObjectBuffer == nullptr) { + hipError_t err = ihipMalloc(&texObjectBuffer, sizeof(__hip_texture), CL_MEM_SVM_FINE_GRAIN_BUFFER); + if (texObjectBuffer == nullptr || err != hipSuccess) { return hipErrorOutOfMemory; } *pTexObject = new (texObjectBuffer) __hip_texture{image, sampler, *pResDesc, *pTexDesc, (pResViewDesc != nullptr) ? *pResViewDesc : hipResourceViewDesc{}};