[Openmp-commits] [PATCH] D102691: [AMDGPU][Libomptarget] Move KernelNameMap to function scope

[Pushpinder Singh via Phabricator via Openmp-commits]

pdhaliwal created this revision.
pdhaliwal added reviewers: JonChesterfield, ronlieb.
Herald added subscribers: kerbowa, t-tye, tpr, dstuttard, yaxunl, nhaehnle, jvesely, kzhuravl.
pdhaliwal requested review of this revision.
Herald added subscribers: openmp-commits, wdng.
Herald added a project: OpenMP.

This patch moves KernelNameMap to function scope where it is
passed on to the called methods.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D102691

Files:
  openmp/libomptarget/plugins/amdgpu/impl/system.cpp


Index: openmp/libomptarget/plugins/amdgpu/impl/system.cpp
===================================================================
--- openmp/libomptarget/plugins/amdgpu/impl/system.cpp
+++ 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));


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D102691.346172.patch
Type: text/x-patch
Size: 3584 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20210518/36cbaae3/attachment-0001.bin>