[libc-commits] [PATCH] D155832: [libc][amdgpu] Accept deadstripped clock_freq global

[Jon Chesterfield via Phabricator via libc-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG095e69404ae8: [libc][amdgpu] Accept deadstripped clock_freq global (authored by JonChesterfield).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D155832/new/

https://reviews.llvm.org/D155832

Files:
  libc/utils/gpu/loader/amdgpu/Loader.cpp


Index: libc/utils/gpu/loader/amdgpu/Loader.cpp
===================================================================
--- libc/utils/gpu/loader/amdgpu/Loader.cpp
+++ libc/utils/gpu/loader/amdgpu/Loader.cpp
@@ -423,32 +423,35 @@
     handle_error(err);
 
   // Obtain the GPU's fixed-frequency clock rate and copy it to the GPU.
-  void *host_clock_freq;
-  if (hsa_status_t err =
-          hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(uint64_t),
-                                       /*flags=*/0, &host_clock_freq))
-    handle_error(err);
-  hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_clock_freq);
-
-  if (hsa_status_t err = hsa_agent_get_info(
-          dev_agent,
-          static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),
-          host_clock_freq))
-    handle_error(err);
-
+  // If the clock_freq symbol is missing, no work to do.
   hsa_executable_symbol_t freq_sym;
-  if (hsa_status_t err = hsa_executable_get_symbol_by_name(
-          executable, "__llvm_libc_clock_freq", &dev_agent, &freq_sym))
-    handle_error(err);
+  if (HSA_STATUS_SUCCESS ==
+      hsa_executable_get_symbol_by_name(executable, "__llvm_libc_clock_freq",
+                                        &dev_agent, &freq_sym)) {
+
+    void *host_clock_freq;
+    if (hsa_status_t err =
+            hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(uint64_t),
+                                         /*flags=*/0, &host_clock_freq))
+      handle_error(err);
+    hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_clock_freq);
 
-  void *freq_addr;
-  if (hsa_status_t err = hsa_executable_symbol_get_info(
-          freq_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &freq_addr))
-    handle_error(err);
+    if (hsa_status_t err =
+            hsa_agent_get_info(dev_agent,
+                               static_cast<hsa_agent_info_t>(
+                                   HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),
+                               host_clock_freq))
+      handle_error(err);
 
-  if (hsa_status_t err = hsa_memcpy(freq_addr, dev_agent, host_clock_freq,
-                                    host_agent, sizeof(uint64_t)))
-    handle_error(err);
+    void *freq_addr;
+    if (hsa_status_t err = hsa_executable_symbol_get_info(
+            freq_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &freq_addr))
+      handle_error(err);
+
+    if (hsa_status_t err = hsa_memcpy(freq_addr, dev_agent, host_clock_freq,
+                                      host_agent, sizeof(uint64_t)))
+      handle_error(err);
+  }
 
   // Obtain a queue with the minimum (power of two) size, used to send commands
   // to the HSA runtime and launch execution on the device.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D155832.542487.patch
Type: text/x-patch
Size: 2722 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/libc-commits/attachments/20230720/70b19177/attachment.bin>