[libc][amdgpu] Accept deadstripped clock_freq global

[JonChesterfield]

Summary

If the clock_freq symbol isn't used, and is removed,
we don't need to abort the loader. Can instead just not set it.

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

auto rpc_alloc = [](uint64_t size, void *data) {
     hsa_amd_agents_allow_access(1, &dev_agent, nullptr, dev_ptr);
     return dev_ptr;
   };
   if (rpc_status_t err = rpc_server_init(device_id, RPC_MAXIMUM_PORT_COUNT,
                                          wavefront_size, rpc_alloc, &tuple))
     handle_error(err);
 
   // Obtain the GPU's fixed-frequency clock rate and copy it to the GPU.
+  // If the clock_freq symbol is missing, no work to do.
+  hsa_executable_symbol_t freq_sym;
+  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);
 
-  if (hsa_status_t err = hsa_agent_get_info(
-          dev_agent,
-          static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),
+    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);
 
-  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);
-
     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.
   uint64_t queue_size;
   if (hsa_status_t err = hsa_agent_get_info(
           dev_agent, HSA_AGENT_INFO_QUEUE_MIN_SIZE, &queue_size))
     handle_error(err);
   hsa_queue_t *queue = nullptr;