HIP resource leaks fix from Jack

Change-Id: I93f3ad7cb94ff1cba1577bd8acc90e826693d12e


[ROCm/hip commit: 778c6626fd]

projects/hip/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<hsa_amd_memory_pool_t*>(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) |

projects/hip/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<ihipFunction_t*> hipFunctionTable;
+};
 
 template <typename MUTEX_TYPE>
 class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>

projects/hip/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<hsa_amd_memory_pool_t*>(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