[libc-commits] [libc] 9553e15 - [libc] Allocate fine-grained memory for the RPC host symbol

[Joseph Huber via libc-commits]

Author: Joseph Huber
Date: 2023-12-01T13:47:33-06:00
New Revision: 9553e156cb840ba4bc040bbfc1f44dc284a97c86

URL: https://github.com/llvm/llvm-project/commit/9553e156cb840ba4bc040bbfc1f44dc284a97c86
DIFF: https://github.com/llvm/llvm-project/commit/9553e156cb840ba4bc040bbfc1f44dc284a97c86.diff

LOG: [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.

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp
index 5c28607b2f290a9..0ff2dce813ed2b8 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(