diff options
author | Jon Chesterfield <jonathanchesterfield@gmail.com> | 2023-07-20 14:23:07 +0100 |
---|---|---|
committer | Jon Chesterfield <jonathanchesterfield@gmail.com> | 2023-07-20 14:23:08 +0100 |
commit | 095e69404ae8feeed32d9dc09debe77702efe7d1 (patch) | |
tree | 2bea9acb965d1677285fb0ca5c7d332de0473a9d /libc | |
parent | 5522e316f91ef490b34ffdf2c21fc8fed8affe5f (diff) | |
download | llvm-095e69404ae8feeed32d9dc09debe77702efe7d1.zip llvm-095e69404ae8feeed32d9dc09debe77702efe7d1.tar.gz llvm-095e69404ae8feeed32d9dc09debe77702efe7d1.tar.bz2 |
[libc][amdgpu] Accept deadstripped clock_freq global
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.
Reviewed By: jhuber6
Differential Revision: https://reviews.llvm.org/D155832
Diffstat (limited to 'libc')
-rw-r--r-- | libc/utils/gpu/loader/amdgpu/Loader.cpp | 49 |
1 files changed, 26 insertions, 23 deletions
diff --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp index c80c4e6..970f9b6 100644 --- a/libc/utils/gpu/loader/amdgpu/Loader.cpp +++ b/libc/utils/gpu/loader/amdgpu/Loader.cpp @@ -423,32 +423,35 @@ int load(int argc, char **argv, char **envp, void *image, size_t size, 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. |