diff --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp
--- a/libc/utils/gpu/loader/amdgpu/Loader.cpp
+++ b/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.