From 9553e156cb840ba4bc040bbfc1f44dc284a97c86 Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Fri, 1 Dec 2023 13:45:10 -0600 Subject: [PATCH] [libc] Allocate fine-grained memory for the RPC host symbol Summary: This pointer has been causing issues. Allocating and reading from coarse memory on the CPU is not guaranteed and varies depending on the kernel version and support. Previously we attempted to pin the memory but this caused unexpected failures. This should be a legal operation and work around the problem as fine-grained memory should be always legal to write to by both sides. --- libc/utils/gpu/loader/amdgpu/Loader.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp index 5c28607b2f29..0ff2dce813ed 100644 --- a/libc/utils/gpu/loader/amdgpu/Loader.cpp +++ b/libc/utils/gpu/loader/amdgpu/Loader.cpp @@ -466,9 +466,10 @@ int load(int argc, char **argv, char **envp, void *image, size_t size, void *rpc_client_host; if (hsa_status_t err = - hsa_amd_memory_pool_allocate(coarsegrained_pool, sizeof(void *), + hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(void *), /*flags=*/0, &rpc_client_host)) handle_error(err); + hsa_amd_agents_allow_access(1, &dev_agent, nullptr, rpc_client_host); void *rpc_client_dev; if (hsa_status_t err = hsa_executable_symbol_get_info(