AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

SWDEV-306940 - Fix hipamd errors/warnings found by CppCheck

Parcourir la source 
Change-Id: Ia93d4f5262ea4269d8031edbfd27033b887e5300
Cette révision appartient à :
Satyanvesh Dittakavi
2021-11-29 17:53:36 +00:00
Parent c792ff7bfc
révision 6e562f4d83
21 fichiers modifiés avec 41 ajouts et 56 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
hipamd/src/cl_d3d10.cpp
-1
Voir le fichier
@@ -310,7 +310,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateImageFromD3D10Resource,
// Check for image support
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
hipamd/src/cl_d3d11.cpp
-1
Voir le fichier
@@ -309,7 +309,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateImageFromD3D11Resource,
// Check for image support
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
hipamd/src/cl_d3d9.cpp
-1
Voir le fichier
@@ -199,7 +199,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromDX9MediaSurfaceKHR,
// Check for image support
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
hipamd/src/cl_gl.cpp
-3
Voir le fichier
@@ -214,7 +214,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromGLTexture,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
@@ -299,7 +298,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromGLTexture2D,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
@@ -379,7 +377,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromGLTexture3D,
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
bool sizePass = false;
for (const auto& it : devices) {
if (it->info().imageSupport_) {
supportPass = true;
hipamd/src/hip_code_object.cpp
+1 -2
Voir le fichier
@@ -624,7 +624,6 @@ hipError_t DynCO::initDynManagedVars(const std::string& managedVar) {
hipError_t DynCO::populateDynGlobalVars() {
amd::ScopedLock lock(dclock_);
std::vector<std::string> var_names;
std::vector<std::string> undef_var_names;
std::string managedVarExt = ".managed";
// For Dynamic Modules there is only one hipFatBinaryDevInfo_
device::Program* dev_program = fb_info_->GetProgram(ihipGetDevice())
@@ -730,7 +729,7 @@ hipError_t StatCO::removeFatBinary(FatBinaryInfo** module) {
if ((*it)->moduleInfo() == module) {
ihipFree((*it)->getManagedVarPtr());
delete *it;
managedVars_.erase(it);
it = managedVars_.erase(it);
} else {
++it;
}
hipamd/src/hip_code_object.hpp
+3 -3
Voir le fichier
@@ -81,12 +81,12 @@ class DynCO : public CodeObject {
amd::Monitor dclock_{"Guards Dynamic Code object", true};
public:
DynCO() : device_id_(ihipGetDevice()) {}
DynCO() : device_id_(ihipGetDevice()), fb_info_(nullptr) {}
virtual ~DynCO();
//LoadsCodeObject and its data
hipError_t loadCodeObject(const char* fname, const void* image=nullptr);
hipModule_t module() { return fb_info_->Module(ihipGetDevice()); };
hipModule_t module() const { return fb_info_->Module(ihipGetDevice()); };
//Gets GlobalVar/Functions from a dynamically loaded code object
hipError_t getDynFunc(hipFunction_t* hfunc, std::string func_name);
@@ -101,7 +101,7 @@ public:
return hipSuccess;
}
// Device ID Check to check if module is launched in the same device it was loaded.
inline void CheckDeviceIdMatch() {
inline void CheckDeviceIdMatch() const {
if (device_id_ != ihipGetDevice()) {
guarantee(false, "Device mismatch from where this module is loaded");
}
hipamd/src/hip_context.cpp
+1 -1
Voir le fichier
@@ -94,9 +94,9 @@ amd::HostQueue* getQueue(hipStream_t stream) {
if (stream == nullptr) {
return getNullStream();
} else {
constexpr bool WaitNullStreamOnly = true;
amd::HostQueue* queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
if (!(reinterpret_cast<hip::Stream*>(stream)->Flags() & hipStreamNonBlocking)) {
constexpr bool WaitNullStreamOnly = true;
iHipWaitActiveStreams(queue, WaitNullStreamOnly);
}
return queue;
hipamd/src/hip_device_runtime.cpp
+3 -1
Voir le fichier
@@ -340,7 +340,9 @@ hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) {
int pciDeviceID = -1;
int pciDomainID = -1;
bool found = false;
if (sscanf (pciBusIdstr, "%04x:%02x:%02x", &pciDomainID, &pciBusID, &pciDeviceID) == 0x3) {
if (sscanf (pciBusIdstr, "%04x:%02x:%02x", reinterpret_cast<unsigned int*>(&pciDomainID),
reinterpret_cast<unsigned int*>(&pciBusID),
reinterpret_cast<unsigned int*>(&pciDeviceID)) == 0x3) {
int count = 0;
ihipDeviceGetCount(&count);
for (cl_int i = 0; i < count; i++) {
hipamd/src/hip_event.hpp
+7 -7
Voir le fichier
@@ -71,8 +71,8 @@ class Event {
std::vector<hipGraphNode_t> nodesPrevToRecorded_;
public:
Event(unsigned int flags)
: flags(flags), lock_("hipEvent_t", true), event_(nullptr), recorded_(false) {
Event(unsigned int flags) : flags(flags), lock_("hipEvent_t", true),
event_(nullptr), recorded_(false), stream_(nullptr) {
// No need to init event_ here as addMarker does that
onCapture_ = false;
device_id_ = hip::getCurrentDevice()->deviceId(); // Created in current device ctx
@@ -87,7 +87,7 @@ class Event {
virtual hipError_t query();
virtual hipError_t synchronize();
hipError_t elapsedTime(Event& stop, float& ms);
hipError_t elapsedTime(Event& eStop, float& ms);
virtual hipError_t streamWaitCommand(amd::Command*& command, amd::HostQueue* queue);
virtual hipError_t enqueueStreamWaitCommand(hipStream_t stream, amd::Command* command);
@@ -107,9 +107,9 @@ class Event {
command.retain();
}
bool isRecorded() { return recorded_; }
bool isRecorded() const { return recorded_; }
amd::Monitor& lock() { return lock_; }
const int deviceId() { return device_id_; }
const int deviceId() const { return device_id_; }
void setDeviceId(int id) { device_id_ = id; }
/// End capture on this event
@@ -123,9 +123,9 @@ class Event {
captureStream_ = stream;
}
/// Get capture status of the graph
bool GetCaptureStatus() { return onCapture_; }
bool GetCaptureStatus() const { return onCapture_; }
/// Get capture stream where event is recorded
hipStream_t GetCaptureStream() { return captureStream_; }
hipStream_t GetCaptureStream() const { return captureStream_; }
/// Set capture stream where event is recorded
void SetCaptureStream(hipStream_t stream) { captureStream_ = stream; }
/// Returns previous captured nodes before event record
hipamd/src/hip_global.cpp
+6 -5
Voir le fichier
@@ -96,7 +96,7 @@ DeviceFunc::~DeviceFunc() {
}
//Abstract functions
Function::Function(std::string name, FatBinaryInfo** modules)
Function::Function(const std::string& name, FatBinaryInfo** modules)
: name_(name), modules_(modules) {
dFunc_.resize(g_devices.size());
}
@@ -167,15 +167,16 @@ hipError_t Function::getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId)
}
//Abstract Vars
Var::Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm,
Var::Var(const std::string& name, DeviceVarKind dVarKind, size_t size, int type, int norm,
FatBinaryInfo** modules) : name_(name), dVarKind_(dVarKind), size_(size),
type_(type), norm_(norm), modules_(modules) {
type_(type), norm_(norm), modules_(modules), managedVarPtr_(nullptr), align_(0) {
dVar_.resize(g_devices.size());
}
Var::Var(std::string name, DeviceVarKind dVarKind, void *pointer, size_t size,
Var::Var(const std::string& name, DeviceVarKind dVarKind, void *pointer, size_t size,
unsigned align, FatBinaryInfo** modules) : name_(name), dVarKind_(dVarKind),
size_(size), modules_(modules), managedVarPtr_(pointer), align_(align) {
size_(size), modules_(modules), managedVarPtr_(pointer), align_(align),
type_(0), norm_(0) {
dVar_.resize(g_devices.size());
}
hipamd/src/hip_global.hpp
+3 -3
Voir le fichier
@@ -57,7 +57,7 @@ private:
//Abstract Structures
class Function {
public:
Function(std::string name, FatBinaryInfo** modules=nullptr);
Function(const std::string& name, FatBinaryInfo** modules=nullptr);
~Function();
//Return DeviceFunc for this this dynamically loaded module
@@ -85,10 +85,10 @@ public:
DVK_Managed
};
Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm,
Var(const std::string& name, DeviceVarKind dVarKind, size_t size, int type, int norm,
FatBinaryInfo** modules = nullptr);
Var(std::string name, DeviceVarKind dVarKind, void *pointer, size_t size, unsigned align,
Var(const std::string& name, DeviceVarKind dVarKind, void *pointer, size_t size, unsigned align,
FatBinaryInfo** modules = nullptr);
~Var();
hipamd/src/hip_graph.cpp
+1 -2
Voir le fichier
@@ -462,10 +462,9 @@ hipError_t capturehipMemcpy(hipStream_t stream, void* dst, const void* src, size
return hipErrorInvalidValue;
}
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
hipGraph_t graph = nullptr;
std::vector<hipGraphNode_t> pDependencies = s->GetLastCapturedNodes();
size_t numDependencies = s->GetLastCapturedNodes().size();
graph = s->GetCaptureGraph();
hipGraph_t graph = s->GetCaptureGraph();
hipError_t status = ihipMemcpy_validate(dst, src, sizeBytes, kind);
if (status != hipSuccess) {
return status;
hipamd/src/hip_graph_internal.hpp
+4 -4
Voir le fichier
@@ -235,9 +235,9 @@ struct ihipGraph {
/// returns all the edges in the graph
std::vector<std::pair<Node, Node>> GetEdges() const;
void GetRunListUtil(Node v, std::unordered_map<Node, bool>& visited,
std::vector<Node>& singleList, std::vector<std::vector<Node>>& parallelList,
std::vector<Node>& singleList, std::vector<std::vector<Node>>& parallelLists,
std::unordered_map<Node, std::vector<Node>>& dependencies);
void GetRunList(std::vector<std::vector<Node>>& parallelList,
void GetRunList(std::vector<std::vector<Node>>& parallelLists,
std::unordered_map<Node, std::vector<Node>>& dependencies);
void LevelOrder(std::vector<Node>& levelOrder);
ihipGraph* clone(std::unordered_map<Node, Node>& clonedNodes) const;
@@ -877,7 +877,7 @@ class hipGraphEventRecordNode : public hipGraphNode {
}
}
void GetParams(hipEvent_t* event) { *event = event_; }
void GetParams(hipEvent_t* event) const { *event = event_; }
hipError_t SetParams(hipEvent_t event) {
event_ = event;
@@ -938,7 +938,7 @@ class hipGraphEventWaitNode : public hipGraphNode {
}
}
void GetParams(hipEvent_t* event) { *event = event_; }
void GetParams(hipEvent_t* event) const { *event = event_; }
hipError_t SetParams(hipEvent_t event) {
event_ = event;
hipamd/src/hip_internal.hpp
+2 -2
Voir le fichier
@@ -306,7 +306,7 @@ namespace hip {
/// Set parent stream
void SetParentStream(hipStream_t parentStream) { parentStream_ = parentStream; }
/// Get parent stream
hipStream_t GetParentStream() { return parentStream_; }
hipStream_t GetParentStream() const { return parentStream_; }
/// Generate ID for stream capture unique over the lifetime of the process
static int GenerateCaptureID() {
static std::atomic<unsigned long long> uid(0);
@@ -408,7 +408,7 @@ namespace hip {
/// Get ROCclr queue associated with hipStream
/// Note: This follows the CUDA spec to sync with default streams
/// and Blocking streams
extern amd::HostQueue* getQueue(hipStream_t s);
extern amd::HostQueue* getQueue(hipStream_t stream);
/// Get default stream associated with the ROCclr context
extern amd::HostQueue* getNullStream(amd::Context&);
/// Get default stream of the thread
hipamd/src/hip_memory.cpp
+3 -8
Voir le fichier
@@ -194,7 +194,6 @@ hipError_t hipSignalExternalSemaphoresAsync(
HIP_RETURN(hipErrorInvalidValue);
}
amd::HostQueue* queue = hip::getQueue(stream);
const amd::Device& device = queue->vdev()->device();
for (unsigned int i = 0; i < numExtSems; i++) {
if (extSemArray[i] != nullptr) {
@@ -224,7 +223,6 @@ hipError_t hipWaitExternalSemaphoresAsync(const hipExternalSemaphore_t* extSemAr
HIP_RETURN(hipErrorInvalidValue);
}
amd::HostQueue* queue = hip::getQueue(stream);
const amd::Device& device = queue->vdev()->device();
for (unsigned int i = 0; i < numExtSems; i++) {
if (extSemArray[i] != nullptr) {
@@ -692,9 +690,8 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
}
const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 };
hipError_t status = hipSuccess;
status = ihipMallocPitch(&pitchedDevPtr->ptr, &pitch, extent.width, extent.height, extent.depth,
CL_MEM_OBJECT_IMAGE3D, &image_format);
hipError_t status = ihipMallocPitch(&pitchedDevPtr->ptr, &pitch, extent.width, extent.height,
extent.depth, CL_MEM_OBJECT_IMAGE3D, &image_format);
if (status == hipSuccess) {
pitchedDevPtr->pitch = pitch;
@@ -1857,7 +1854,6 @@ inline hipError_t ihipMemcpyCmdEnqueue(amd::Command* command, bool isAsync = fal
}
hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool isAsync = false) {
amd::Command* command;
hipError_t status;
if (pCopy == nullptr || !hip::isValid(stream)) {
return hipErrorInvalidValue;
@@ -1904,6 +1900,7 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool
pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch,
pCopy->dstPitch * pCopy->dstHeight);
} else {
amd::Command* command;
status = ihipGetMemcpyParam3DCommand(command, pCopy, hip::getQueue(stream));
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
@@ -2298,7 +2295,6 @@ hipError_t ihipMemsetCommand(std::vector<amd::Command*>& commands, void* dst, in
amd::Memory* memory = getMemoryObject(dst, offset);
size_t n_head_bytes = 0;
size_t n_tail_bytes = 0;
int64_t value64 = 0;
amd::Command* command;
hip_error = packFillMemoryCommand(command, memory, offset, value, valueSize, sizeBytes,
@@ -2569,7 +2565,6 @@ hipError_t hipIpcOpenMemHandle(void** dev_ptr, hipIpcMemHandle_t handle, unsigne
hipError_t hipIpcCloseMemHandle(void* dev_ptr) {
HIP_INIT_API(hipIpcCloseMemHandle, dev_ptr);
size_t offset = 0;
amd::Device* device = nullptr;
amd::Memory* amd_mem_obj = nullptr;
hipamd/src/hip_module.cpp
+2 -3
Voir le fichier
@@ -324,10 +324,9 @@ hipError_t ihipLaunchKernelCommand(amd::Command*& command, hipFunction_t f,
size_t localWorkSize[3] = {blockDimX, blockDimY, blockDimZ};
amd::NDRangeContainer ndrange(3, globalWorkOffset, globalWorkSize, localWorkSize);
amd::Command::EventWaitList waitList;
bool profileNDRange = false;
address kernargs = nullptr;
profileNDRange = (startEvent != nullptr || stopEvent != nullptr);
bool profileNDRange = (startEvent != nullptr || stopEvent != nullptr);
// Flag set to 1 signifies that kernel can be launched in anyorder
if (flags & hipExtAnyOrderLaunch) {
@@ -553,10 +552,10 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
hipError_t result = hipErrorUnknown;
uint64_t allGridSize = 0;
uint32_t blockDims = 0;
std::vector<const amd::Device*> mgpu_list(numDevices);
for (int i = 0; i < numDevices; ++i) {
uint32_t blockDims = 0;
const hipLaunchParams& launch = launchParamsList[i];
blockDims = launch.blockDim.x * launch.blockDim.y * launch.blockDim.z;
allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z *
hipamd/src/hip_peer.cpp
-1
Voir le fichier
@@ -161,7 +161,6 @@ hipError_t hipDeviceGetP2PAttribute(int* value, hipDeviceP2PAttr attr,
default : {
LogPrintfError("Invalid attribute attr: %d ", attr);
HIP_RETURN(hipErrorInvalidValue);
break;
}
}
hipamd/src/hip_platform.cpp
+3 -5
Voir le fichier
@@ -292,12 +292,9 @@ hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memor
{
HIP_INIT();
amd::Program* program = nullptr;
device::Program* dev_program = nullptr;
/* Get Device Program pointer*/
program = as_amd(reinterpret_cast<cl_program>(hmod));
dev_program = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
device::Program* dev_program = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
if (dev_program == nullptr) {
LogPrintfError("Cannot get Device Function for module: 0x%x \n", hmod);
@@ -761,6 +758,7 @@ hipError_t PlatformState::loadModule(hipModule_t *module, const char* fname, con
assert(*module != nullptr);
if (dynCO_map_.find(*module) != dynCO_map_.end()) {
delete dynCo;
return hipErrorAlreadyMapped;
}
dynCO_map_.insert(std::make_pair(*module, dynCo));
hipamd/src/hip_platform.hpp
+1 -1
Voir le fichier
@@ -28,7 +28,7 @@ namespace hip_impl {
hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
int* maxBlocksPerCU, int* numBlocksPerGrid, int* bestBlockSize,
const amd::Device& device, hipFunction_t func, int blockSize,
const amd::Device& device, hipFunction_t func, int inputBlockSize,
size_t dynamicSMemSize, bool bCalcPotentialBlkSz);
} /* namespace hip_impl*/
hipamd/src/hip_rtc.cpp
-1
Voir le fichier
@@ -127,7 +127,6 @@ static std::string getValueOf(const std::string& option) {
}
static void transformOptions(std::vector<std::string>& options, amd::Program* program) {
std::vector<const char*> t_option;
for (auto& i : options) {
// This functionality is no longer supported - just consume at the moment and remove this in
// couple of releases
hipamd/src/hip_stream.cpp
+1 -1
Voir le fichier
@@ -426,7 +426,7 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
}
void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_data) {
CallbackData* data = (CallbackData*)user_data;
CallbackData* data = reinterpret_cast<CallbackData*>(user_data);
int offset = data->previous_read_index % IPC_SIGNALS_PER_EVENT;
while (data->shmem->read_index < data->previous_read_index + IPC_SIGNALS_PER_EVENT &&
data->shmem->signal[offset] != 0) {