[libc-commits] [libc] 9937952 - [libc][Docs] Update RPC server example using CUDA after changes

[Joseph Huber via libc-commits]

Author: Joseph Huber
Date: 2024-03-29T13:13:20-05:00
New Revision: 99379522d00d0720bcda5067d4de2bfb0c279f15

URL: https://github.com/llvm/llvm-project/commit/99379522d00d0720bcda5067d4de2bfb0c279f15
DIFF: https://github.com/llvm/llvm-project/commit/99379522d00d0720bcda5067d4de2bfb0c279f15.diff

LOG: [libc][Docs] Update RPC server example using CUDA after changes

Summary:
This has changed, so update it to match the new interface.

Added: 
    

Modified: 
    libc/docs/gpu/rpc.rst

Removed: 
    


################################################################################
diff  --git a/libc/docs/gpu/rpc.rst b/libc/docs/gpu/rpc.rst
index 9d6d8099db951c..e13a377f305c06 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 \
        -I<install-path>include -L<install-path>/lib -lllvmlibc_rpc_server \
        -O3 -foffload-lto -o hello
   $> ./hello