AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

SWDEV-301667 - Fix build warnings

Kaynağa Gözat 
Change-Id: Ice23b3d1a19a7958ca5c3caff34db2934b361a98
Bu işleme şunda yer alıyor:
Saleel Kudchadker
2022-03-02 11:46:56 -08:00
ebeveyn 545cfaf428
işleme 086aee0fcb
8 değiştirilmiş dosya ile 76 ekleme ve 54 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
hipamd/src/hip_code_object.cpp
+16 -11
Dosyayı Görüntüle
@@ -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<std::string> 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<address>(dvar->device_ptr()), var->getManagedVarPtr(),
dvar->size(), hipMemcpyHostToDevice, *queue);
err = ihipMemcpy(reinterpret_cast<address>(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
hipamd/src/hip_device_runtime.cpp
+11 -13
Dosyayı Görüntüle
@@ -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<unsigned int*>(&pciBusID),
reinterpret_cast<unsigned int*>(&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,
hipamd/src/hip_global.cpp
+1 -1
Dosyayı Görüntüle
@@ -62,7 +62,7 @@ DeviceVar::~DeviceVar() {
if (shadowVptr != nullptr) {
textureReference* texRef = reinterpret_cast<textureReference*>(shadowVptr);
ihipUnbindTexture(texRef);
hipError_t err = ihipUnbindTexture(texRef);
delete texRef;
shadowVptr = nullptr;
}
hipamd/src/hip_module.cpp
+8 -5
Dosyayı Görüntüle
@@ -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<const amd::Device*> 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);
}
hipamd/src/hip_platform.cpp
+31 -18
Dosyayı Görüntüle
@@ -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<hip::FatBinaryInfo**>(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<void*>(function_address), deviceId);
hipError_t hip_error =
PlatformState::instance().getStatFunc(&func,
reinterpret_cast<void*>(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<void*>(function_address),
hipError_t err = hipLaunchCooperativeKernel(reinterpret_cast<void*>(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;
hipamd/src/hip_stream.cpp
+4 -2
Dosyayı Görüntüle
@@ -62,7 +62,8 @@ hipError_t Stream::EndCapture() {
}
for (auto stream : parallelCaptureStreams_) {
hip::Stream* s = reinterpret_cast<hip::Stream*>(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;
}
hipamd/src/hip_surface.cpp
+3 -2
Dosyayı Görüntüle
@@ -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};
hipamd/src/hip_texture.cpp
+2 -2
Dosyayı Görüntüle
@@ -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{}};