From 778c6626fdf5791b2442fe47079584075d0837fa Mon Sep 17 00:00:00 2001 From: pensun Date: Mon, 5 Dec 2016 20:21:33 -0600 Subject: [PATCH] HIP resource leaks fix from Jack Change-Id: I93f3ad7cb94ff1cba1577bd8acc90e826693d12e --- src/hip_hcc.cpp | 10 ++------- src/hip_hcc.h | 34 ++++++++++++++++++++---------- src/hip_module.cpp | 52 +++++++++++++++++++++++++++++++++++++++++++--- 3 files changed, 74 insertions(+), 22 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index abd260762f..84794b02fb 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -401,17 +401,11 @@ void ihipStream_t::launchModuleKernel( uint32_t gridDimY, uint32_t gridDimZ, uint32_t groupSegmentSize, - uint32_t privateSegmentSize, + uint32_t privateSegmentSize, void *kernarg, size_t kernSize, uint64_t kernel){ hsa_status_t status; - void *kern; - - hsa_amd_memory_pool_t *pool = reinterpret_cast(av.get_hsa_kernarg_region()); - status = hsa_amd_memory_pool_allocate(*pool, kernSize, 0, &kern); - status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)av.get_hsa_agent(), 0, kern); - memcpy(kern, kernarg, kernSize); hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue(); const uint32_t queue_mask = Queue->size-1; uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue); @@ -426,7 +420,7 @@ void ihipStream_t::launchModuleKernel( dispatch_packet->grid_size_z = blockDimZ * gridDimZ; dispatch_packet->group_segment_size = groupSegmentSize; dispatch_packet->private_segment_size = privateSegmentSize; - dispatch_packet->kernarg_address = kern; + dispatch_packet->kernarg_address = kernarg; dispatch_packet->kernel_object = kernel; uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | diff --git a/src/hip_hcc.h b/src/hip_hcc.h index b01d41be14..3074a4f121 100644 --- a/src/hip_hcc.h +++ b/src/hip_hcc.h @@ -367,17 +367,6 @@ struct LockedBase { MUTEX_TYPE _mutex; }; - -class ihipModule_t{ -public: - hsa_executable_t executable; - hsa_code_object_t object; - std::string fileName; - void *ptr; - size_t size; -}; - - class ihipFunction_t{ public: ihipFunction_t(const char *name) { @@ -399,6 +388,29 @@ public: uint64_t _kernel; }; +class ihipModule_t { +public: + hsa_executable_t executable; + hsa_code_object_t object; + std::string fileName; + void *ptr; + size_t size; + + ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0), hipFunctionTable() {} + ~ihipModule_t() { + for (int i = 0; i < hipFunctionTable.size(); ++i) { + ihipFunction_t *func = hipFunctionTable[i]; + delete func; + } + hipFunctionTable.clear(); + } + + void registerFunction(ihipFunction_t* func) { + hipFunctionTable.push_back(func); + } +private: + std::vector hipFunctionTable; +}; template class ihipStreamCriticalBase_t : public LockedBase diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 606d99f2fd..d58a476f77 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -126,7 +126,6 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ } else { - *module = new ihipModule_t; size_t size = std::string::size_type(in.tellg()); void *p = NULL; hsa_agent_t agent = currentDevice->_hsaAgent; @@ -172,6 +171,11 @@ hipError_t hipModuleUnload(hipModule_t hmod){ ret = hipErrorInvalidValue; } status = hsa_code_object_destroy(hmod->object); + if(status != HSA_STATUS_SUCCESS) + { + ret = hipErrorInvalidValue; + } + status = hsa_memory_free(hmod->ptr); if(status != HSA_STATUS_SUCCESS) { ret = hipErrorInvalidValue; @@ -193,6 +197,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch }else{ *func = new ihipFunction_t(name); + hmod->registerFunction(*func); int deviceId = ctx->getDevice()->_deviceId; ihipDevice_t *currentDevice = ihipGetDevice(deviceId); hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent; @@ -268,11 +273,17 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &groupSegmentSize); + if(status != HSA_STATUS_SUCCESS){ + return ihipLogStatus(hipErrorNotFound); + } uint32_t privateSegmentSize; status = hsa_executable_symbol_get_info(f->_kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &privateSegmentSize); + if(status != HSA_STATUS_SUCCESS){ + return ihipLogStatus(hipErrorNotFound); + } privateSegmentSize += sharedMemBytes; @@ -300,7 +311,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, aql.group_segment_size = groupSegmentSize; aql.private_segment_size = privateSegmentSize; aql.kernel_object = f->_kernel; - aql.setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; + aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | @@ -315,13 +326,32 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, hsa_signal_t signal; status = hsa_signal_create(1, 0, NULL, &signal); + if(status != HSA_STATUS_SUCCESS){ + return ihipLogStatus(hipErrorNotFound); + } + + /* + Allocate kernarg + */ + void *kern = nullptr; + + hsa_amd_memory_pool_t *pool = reinterpret_cast(lp.av->get_hsa_kernarg_region()); + status = hsa_amd_memory_pool_allocate(*pool, kernArgSize, 0, &kern); + if(status != HSA_STATUS_SUCCESS){ + return ihipLogStatus(hipErrorNotFound); + } + status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)lp.av->get_hsa_agent(), 0, kern); + if(status != HSA_STATUS_SUCCESS){ + return ihipLogStatus(hipErrorNotFound); + } + memcpy(kern, config[1], kernArgSize); /* Launch AQL packet */ hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ, - gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, config[1], kernArgSize, f->_kernel); + gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_kernel); /* @@ -330,6 +360,22 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + /* + Destroy kernarg + */ + status = hsa_amd_memory_pool_free(kern); + if(status != HSA_STATUS_SUCCESS){ + return ihipLogStatus(hipErrorNotFound); + } + + /* + Destroy the signal + */ + status = hsa_signal_destroy(signal); + if(status != HSA_STATUS_SUCCESS){ + return ihipLogStatus(hipErrorNotFound); + } + #endif // USE_DISPATCH_HSA_KERNEL