diff --git a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
--- a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
@@ -132,7 +132,6 @@
 
 std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
 
-std::map<std::string, std::string> KernelNameMap;
 std::vector<std::map<std::string, atl_kernel_info_t>> KernelInfoTable;
 std::vector<std::map<std::string, atl_symbol_info_t>> SymbolInfoTable;
 
@@ -833,8 +832,9 @@
 }
 } // namespace
 
-static hsa_status_t get_code_object_custom_metadata(void *binary,
-                                                    size_t binSize, int gpu) {
+static hsa_status_t get_code_object_custom_metadata(
+    std::map<std::string, std::string> &KernelNameMap, void *binary,
+    size_t binSize, int gpu) {
   // parse code object with different keys from v2
   // also, the kernel name is not the same as the symbol name -- so a
   // symbol->name map is needed
@@ -1005,10 +1005,26 @@
   return HSA_STATUS_SUCCESS;
 }
 
-static hsa_status_t populate_InfoTables(hsa_executable_t executable,
-                                        hsa_executable_symbol_t symbol,
-                                        void *data) {
-  int gpu = *static_cast<int *>(data);
+struct PopulateInfoTablesContext {
+  PopulateInfoTablesContext(std::map<std::string, std::string> &KernelNameMap,
+                            int Device)
+      : KernelNameMap(KernelNameMap), DeviceId(Device) {}
+
+  PopulateInfoTablesContext() = delete;
+  PopulateInfoTablesContext(const PopulateInfoTablesContext &) = delete;
+  ~PopulateInfoTablesContext() = default;
+
+  std::map<std::string, std::string> &KernelNameMap;
+  int DeviceId;
+};
+
+static hsa_status_t PopulateInfoTables(hsa_executable_t executable,
+                                       hsa_executable_symbol_t symbol,
+                                       void *data) {
+  PopulateInfoTablesContext *Ctx =
+      static_cast<PopulateInfoTablesContext *>(data);
+  int gpu = Ctx->DeviceId;
+  auto &KernelNameMap = Ctx->KernelNameMap;
   hsa_symbol_kind_t type;
 
   uint32_t name_length;
@@ -1162,6 +1178,7 @@
     void *cb_state, std::vector<hsa_executable_t> &HSAExecutables) {
   hsa_status_t err;
   int gpu = place.device_id;
+  std::map<std::string, std::string> KernelNameMap;
   assert(gpu >= 0);
 
   DEBUG_PRINT("Trying to load module to GPU-%d\n", gpu);
@@ -1195,7 +1212,8 @@
       // Some metadata info is not available through ROCr API, so use custom
       // code object metadata parsing to collect such metadata info
 
-      err = get_code_object_custom_metadata(module_bytes, module_size, gpu);
+      err = get_code_object_custom_metadata(KernelNameMap, module_bytes,
+                                            module_size, gpu);
       if (err != HSA_STATUS_SUCCESS) {
         DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
                     "Getting custom code object metadata",
@@ -1249,8 +1267,9 @@
       exit(1);
     }
 
-    err = hsa_executable_iterate_symbols(executable, populate_InfoTables,
-                                         static_cast<void *>(&gpu));
+    PopulateInfoTablesContext Ctx(KernelNameMap, gpu);
+    err = hsa_executable_iterate_symbols(executable, PopulateInfoTables,
+                                         static_cast<void *>(&Ctx));
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Iterating over symbols for execuatable", get_error_string(err));