aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorJoseph Huber <huberjn@outlook.com>2024-03-29 13:13:03 -0500
committerJoseph Huber <huberjn@outlook.com>2024-03-29 13:13:20 -0500
commit99379522d00d0720bcda5067d4de2bfb0c279f15 (patch)
tree270e8f16b4da7ba28a26fd69e06b211579f52a33
parenta1a8bb1d3ae9a535526aba9514e87f76e2d040fa (diff)
downloadllvm-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.rst15
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