Fix memory leak introduced by previous change to Agent_global.

Make Agent_global manage the lifetime of the name string


[ROCm/clr commit: fa564a5345]

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -2506,6 +2506,40 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con
 hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func);
 
 struct Agent_global {
+
+    Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
+    Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) 
+      : name(nullptr), address(address), byte_cnt(byte_cnt) {
+      if (name)
+        this->name = strdup(name);
+    }
+
+    Agent_global& operator=(Agent_global&& t) {
+      if (this == &t) return *this;
+
+      if (name) free(name);
+      name = t.name;
+      address = t.address;
+      byte_cnt = t.byte_cnt;
+
+      t.name = nullptr;
+      t.address = nullptr;
+      t.byte_cnt = 0;
+
+      return *this;
+    }
+
+    Agent_global(Agent_global&& t) 
+      : name(nullptr), address(nullptr), byte_cnt(0) {
+      *this = std::move(t);
+    }
+
+    // not needed, delete them to prevent bugs
+    Agent_global(const Agent_global&) = delete;
+    Agent_global& operator=(Agent_global& t) = delete;
+
+    ~Agent_global() { if (name) free(name); }
+    
     char* name;
     hipDeviceptr_t address;
     uint32_t byte_cnt;

projects/clr/hipamd/src/hip_module.cpp

@@ -341,10 +341,8 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag
     hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t);
 
     if (t == HSA_SYMBOL_KIND_VARIABLE) {
-        Agent_global tmp = {nullptr, address(x), size(x)};
-        tmp.name = (char *) malloc(sizeof(name(x).c_str()));
-        strcpy(tmp.name, name(x).c_str());
-        static_cast<Container*>(out)->push_back(tmp);
+        Agent_global tmp(name(x).c_str(), address(x), size(x));
+        static_cast<Container*>(out)->push_back(std::move(tmp));
 
         track(static_cast<Container*>(out)->back(),agent);
     }