diff --git a/hipamd/src/cl_d3d10.cpp b/hipamd/src/cl_d3d10.cpp index 0402f6e051..692b26cb7e 100644 --- a/hipamd/src/cl_d3d10.cpp +++ b/hipamd/src/cl_d3d10.cpp @@ -310,7 +310,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateImageFromD3D10Resource, // Check for image support const std::vector& devices = as_amd(context)->devices(); bool supportPass = false; - bool sizePass = false; for (const auto& it : devices) { if (it->info().imageSupport_) { supportPass = true; diff --git a/hipamd/src/cl_d3d11.cpp b/hipamd/src/cl_d3d11.cpp index dfc2408d45..9ad60cd2ab 100644 --- a/hipamd/src/cl_d3d11.cpp +++ b/hipamd/src/cl_d3d11.cpp @@ -309,7 +309,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateImageFromD3D11Resource, // Check for image support const std::vector& devices = as_amd(context)->devices(); bool supportPass = false; - bool sizePass = false; for (const auto& it : devices) { if (it->info().imageSupport_) { supportPass = true; diff --git a/hipamd/src/cl_d3d9.cpp b/hipamd/src/cl_d3d9.cpp index 956dbfb359..1b78b68054 100644 --- a/hipamd/src/cl_d3d9.cpp +++ b/hipamd/src/cl_d3d9.cpp @@ -199,7 +199,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromDX9MediaSurfaceKHR, // Check for image support const std::vector& devices = as_amd(context)->devices(); bool supportPass = false; - bool sizePass = false; for (const auto& it : devices) { if (it->info().imageSupport_) { supportPass = true; diff --git a/hipamd/src/cl_gl.cpp b/hipamd/src/cl_gl.cpp index 6f232ecaf0..b3730f61da 100644 --- a/hipamd/src/cl_gl.cpp +++ b/hipamd/src/cl_gl.cpp @@ -214,7 +214,6 @@ RUNTIME_ENTRY_RET(cl_mem, clCreateFromGLTexture, const std::vector& 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& 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& devices = as_amd(context)->devices(); bool supportPass = false; - bool sizePass = false; for (const auto& it : devices) { if (it->info().imageSupport_) { supportPass = true; diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index c719a68024..843ea24591 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -624,7 +624,6 @@ hipError_t DynCO::initDynManagedVars(const std::string& managedVar) { hipError_t DynCO::populateDynGlobalVars() { amd::ScopedLock lock(dclock_); std::vector var_names; - std::vector 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; } diff --git a/hipamd/src/hip_code_object.hpp b/hipamd/src/hip_code_object.hpp index 009144ea89..112bce8dd9 100644 --- a/hipamd/src/hip_code_object.hpp +++ b/hipamd/src/hip_code_object.hpp @@ -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"); } diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp index 8c1dfc2933..509cca908d 100644 --- a/hipamd/src/hip_context.cpp +++ b/hipamd/src/hip_context.cpp @@ -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(stream)->asHostQueue(); if (!(reinterpret_cast(stream)->Flags() & hipStreamNonBlocking)) { + constexpr bool WaitNullStreamOnly = true; iHipWaitActiveStreams(queue, WaitNullStreamOnly); } return queue; diff --git a/hipamd/src/hip_device_runtime.cpp b/hipamd/src/hip_device_runtime.cpp index e6898fd183..0e9733d94e 100644 --- a/hipamd/src/hip_device_runtime.cpp +++ b/hipamd/src/hip_device_runtime.cpp @@ -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(&pciDomainID), + reinterpret_cast(&pciBusID), + reinterpret_cast(&pciDeviceID)) == 0x3) { int count = 0; ihipDeviceGetCount(&count); for (cl_int i = 0; i < count; i++) { diff --git a/hipamd/src/hip_event.hpp b/hipamd/src/hip_event.hpp index f727aff4a8..831083d001 100644 --- a/hipamd/src/hip_event.hpp +++ b/hipamd/src/hip_event.hpp @@ -71,8 +71,8 @@ class Event { std::vector 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 diff --git a/hipamd/src/hip_global.cpp b/hipamd/src/hip_global.cpp index b6c05523dd..0d9c2b44ab 100644 --- a/hipamd/src/hip_global.cpp +++ b/hipamd/src/hip_global.cpp @@ -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()); } diff --git a/hipamd/src/hip_global.hpp b/hipamd/src/hip_global.hpp index 6fd9346691..6c964b625d 100644 --- a/hipamd/src/hip_global.hpp +++ b/hipamd/src/hip_global.hpp @@ -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(); diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index d594893831..79a29e536e 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -462,10 +462,9 @@ hipError_t capturehipMemcpy(hipStream_t stream, void* dst, const void* src, size return hipErrorInvalidValue; } hip::Stream* s = reinterpret_cast(stream); - hipGraph_t graph = nullptr; std::vector 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; diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 7c7da322e3..07040d8f8c 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -235,9 +235,9 @@ struct ihipGraph { /// returns all the edges in the graph std::vector> GetEdges() const; void GetRunListUtil(Node v, std::unordered_map& visited, - std::vector& singleList, std::vector>& parallelList, + std::vector& singleList, std::vector>& parallelLists, std::unordered_map>& dependencies); - void GetRunList(std::vector>& parallelList, + void GetRunList(std::vector>& parallelLists, std::unordered_map>& dependencies); void LevelOrder(std::vector& levelOrder); ihipGraph* clone(std::unordered_map& 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; diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index 0bcab97f8e..519e1c206b 100644 --- a/hipamd/src/hip_internal.hpp +++ b/hipamd/src/hip_internal.hpp @@ -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 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 diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 91903a7d7d..1921416f08 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -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& 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; diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 6de77c328d..a4f1fbfc70 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -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 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 * diff --git a/hipamd/src/hip_peer.cpp b/hipamd/src/hip_peer.cpp index 09e90d5a6a..c5e9c2f5f5 100644 --- a/hipamd/src/hip_peer.cpp +++ b/hipamd/src/hip_peer.cpp @@ -161,7 +161,6 @@ hipError_t hipDeviceGetP2PAttribute(int* value, hipDeviceP2PAttr attr, default : { LogPrintfError("Invalid attribute attr: %d ", attr); HIP_RETURN(hipErrorInvalidValue); - break; } } diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp index 12e9aa21d0..bcae85601c 100644 --- a/hipamd/src/hip_platform.cpp +++ b/hipamd/src/hip_platform.cpp @@ -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(hmod)); - dev_program = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); + amd::Program* program = as_amd(reinterpret_cast(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)); diff --git a/hipamd/src/hip_platform.hpp b/hipamd/src/hip_platform.hpp index 6b7cead2fd..4bed5f2ef1 100644 --- a/hipamd/src/hip_platform.hpp +++ b/hipamd/src/hip_platform.hpp @@ -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*/ diff --git a/hipamd/src/hip_rtc.cpp b/hipamd/src/hip_rtc.cpp index f13a7c718c..60daa26492 100644 --- a/hipamd/src/hip_rtc.cpp +++ b/hipamd/src/hip_rtc.cpp @@ -127,7 +127,6 @@ static std::string getValueOf(const std::string& option) { } static void transformOptions(std::vector& options, amd::Program* program) { - std::vector t_option; for (auto& i : options) { // This functionality is no longer supported - just consume at the moment and remove this in // couple of releases diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index 9c8ee495f0..d3720526b8 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -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(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) {