v2: Fixed hipModule memory management

1. Changed test to assert for same hipFunction values
2. Added better memory management for hipModule

Change-Id: I10d7aef13c215a2211e262f3c79017f26a17d9a7

src/hip_hcc.h

@@ -343,7 +343,7 @@ public:
   std::string fileName;
   void *ptr;
   size_t size;
-
+  std::list<hipFunction_t> funcTrack;
   ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0) {}
 };
 

src/hip_module.cpp

@@ -40,11 +40,9 @@ struct ihipModuleSymbol_t{
     uint64_t    _object;             // The kernel object.
     uint32_t    _groupSegmentSize;
     uint32_t    _privateSegmentSize;
-    char        _name[64];       // TODO - review for performance cost.  Name is just used for debug.
+    std::string        _name;       // TODO - review for performance cost.  Name is just used for debug.
 };
 
-std::list<hipFunction_t> hipFuncTracker;
-
 template <>
 std::string ToString(hipFunction_t v)
 {
@@ -236,6 +234,13 @@ hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char
         ret = hipErrorInvalidContext;
 
     }else{
+        std::string str(name);
+        for(std::list<hipFunction_t>::iterator f = hmod->funcTrack.begin(); f != hmod->funcTrack.end(); ++f) {
+          if((*f)->_name == str) {
+            *func = *f;
+          }
+          return ret;
+        }
         ihipModuleSymbol_t *sym = new ihipModuleSymbol_t;
         int deviceId = ctx->getDevice()->_deviceId;
         ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
@@ -263,9 +268,9 @@ hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char
                                     &sym->_privateSegmentSize);
         CHECK_HSA(status, hipErrorNotFound);
 
-        strncpy(sym->_name, name, sizeof(sym->_name));
+        sym->_name = name;
         *func = sym;
-        hipFuncTracker.push_back(*func);
+        hmod->funcTrack.push_back(*func);
     }
     return ihipLogStatus(ret);
 }
@@ -321,7 +326,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         */
         grid_launch_parm lp;
         lp.dynamic_group_mem_bytes = sharedMemBytes;  // TODO - this should be part of preLaunchKernel.
-        hStream = ihipPreLaunchKernel(hStream, dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), &lp, f->_name);
+        hStream = ihipPreLaunchKernel(hStream, dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), &lp, f->_name.c_str());
 
 
         hsa_kernel_dispatch_packet_t aql;
@@ -355,7 +360,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, nullptr/*completion_future*/);
 
 
-        ihipPostLaunchKernel(f->_name, hStream, lp);
+        ihipPostLaunchKernel(f->_name.c_str(), hStream, lp);
     }
 
     return ihipLogStatus(ret);

tests/src/runtimeApi/module/hipModule.cpp

@@ -46,7 +46,6 @@ int main(){
   for(uint32_t i=0;i<LEN;i++){
     A[i] = i*1.0f;
     B[i] = 0.0f;
-    std::cout<<A[i] << " "<<B[i]<<std::endl;
   }
 
   HIPCHECK(hipMalloc((void**)&Ad, SIZE));
@@ -58,41 +57,41 @@ int main(){
   hipFunction_t Function;
   HIPCHECK(hipModuleLoad(&Module, fileName));
   HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
+  hipFunction_t f;
+  HIPCHECK(hipModuleGetFunction(&f, Module, kernel_name));
+  assert(f == Function);
   hipStream_t stream;
   HIPCHECK(hipStreamCreate(&stream));
   void *args[2] = {&Ad, &Bd};
-  std::cout<<Function<<std::endl;
 
-    std::vector<void*>argBuffer(5);
-    memcpy(&argBuffer[3], &Ad, sizeof(void*));
-    memcpy(&argBuffer[4], &Bd, sizeof(void*));
+  std::vector<void*>argBuffer(5);
+  memcpy(&argBuffer[3], &Ad, sizeof(void*));
+  memcpy(&argBuffer[4], &Bd, sizeof(void*));
 
-    size_t size = argBuffer.size()*sizeof(void*);
+  size_t size = argBuffer.size()*sizeof(void*);
 
-    void *config[] = {
+  void *config[] = {
       HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
       HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
       HIP_LAUNCH_PARAM_END
-    };
+  };
 
-    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config);
+  hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config);
 
-    HIPCHECK(hipStreamDestroy(stream));
+  HIPCHECK(hipStreamDestroy(stream));
 
   HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
 
   for(uint32_t i=0;i<LEN;i++){
-  std::cout<<A[i]<<" - "<<B[i]<<std::endl;
+    assert(A[i] == B[i]);
   }
+
   std::vector<hipFunction_t> vec(1024*1024*64);
   for(unsigned i=0;i<1024*1024*64;i++) {
     hipFunction_t func;
     hipModuleGetFunction(&func, Module, kernel_name);
     vec[i] = func;
   }
-
-  std::cout<<"Starting sleep"<<std::endl;
-  std::this_thread::sleep_for(std::chrono::seconds(10));
-  std::cout<<"Done sleeping"<<std::endl;
+  passed();
   return 0;
 }