From 99379522d00d0720bcda5067d4de2bfb0c279f15 Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Fri, 29 Mar 2024 13:13:03 -0500 Subject: [libc][Docs] Update RPC server example using CUDA after changes Summary: This has changed, so update it to match the new interface. --- libc/docs/gpu/rpc.rst | 15 ++++----------- 1 file changed, 4 insertions(+), 11 deletions(-) diff --git a/libc/docs/gpu/rpc.rst b/libc/docs/gpu/rpc.rst index 9d6d809..e13a377 100644 --- a/libc/docs/gpu/rpc.rst +++ b/libc/docs/gpu/rpc.rst @@ -251,14 +251,10 @@ but the following example shows how it can be used by a standard user. __global__ void hello() { puts("Hello world!"); } int main() { - int device = 0; - // Initialize the RPC server to run on a single device. - if (rpc_status_t err = rpc_init(/*num_device=*/1)) - handle_error(err); - // Initialize the RPC server to run on the given device. + rpc_device_t device; if (rpc_status_t err = - rpc_server_init(device, RPC_MAXIMUM_PORT_COUNT, + rpc_server_init(&device, RPC_MAXIMUM_PORT_COUNT, /*warp_size=*/32, alloc_host, /*data=*/nullptr)) handle_error(err); @@ -277,6 +273,7 @@ but the following example shows how it can be used by a standard user. hello<<<1, 1, 0, stream>>>(); // While the kernel is executing, check the RPC server for work to do. + // Requires non-blocking CUDA kernels but avoids a separate thread. while (cudaStreamQuery(stream) == cudaErrorNotReady) if (rpc_status_t err = rpc_handle_server(device)) handle_error(err); @@ -286,10 +283,6 @@ but the following example shows how it can be used by a standard user. rpc_server_shutdown(device, free_host, /*data=*/nullptr)) handle_error(err); - // Shut down the entire RPC server interface. - if (rpc_status_t err = rpc_shutdown()) - handle_error(err); - return EXIT_SUCCESS; } @@ -300,7 +293,7 @@ associated with relocatable device code linking. .. code-block:: sh - $> clang++ -x cuda rpc.cpp --offload-arch=native -fgpu-rdc -lcudart -lcgpu \ + $> clang++ -x cuda rpc.cpp --offload-arch=native -fgpu-rdc -lcudart -lcgpu-nvptx \ -Iinclude -L/lib -lllvmlibc_rpc_server \ -O3 -foffload-lto -o hello $> ./hello -- cgit v1.1