diff options
author | Joseph Huber <huberjn@outlook.com> | 2024-03-29 13:13:03 -0500 |
---|---|---|
committer | Joseph Huber <huberjn@outlook.com> | 2024-03-29 13:13:20 -0500 |
commit | 99379522d00d0720bcda5067d4de2bfb0c279f15 (patch) | |
tree | 270e8f16b4da7ba28a26fd69e06b211579f52a33 | |
parent | a1a8bb1d3ae9a535526aba9514e87f76e2d040fa (diff) | |
download | llvm-99379522d00d0720bcda5067d4de2bfb0c279f15.zip llvm-99379522d00d0720bcda5067d4de2bfb0c279f15.tar.gz llvm-99379522d00d0720bcda5067d4de2bfb0c279f15.tar.bz2 |
[libc][Docs] Update RPC server example using CUDA after changes
Summary:
This has changed, so update it to match the new interface.
-rw-r--r-- | libc/docs/gpu/rpc.rst | 15 |
1 files 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 \ -I<install-path>include -L<install-path>/lib -lllvmlibc_rpc_server \ -O3 -foffload-lto -o hello $> ./hello |