[OpenMP] Add documentation on using the libc in OpenMP

This points users to the `libc` documentation and explains the basics of
how it's used inside the runtime.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D155318

openmp/docs/SupportAndFAQ.rst

@@ -433,3 +433,21 @@ Clang compiler and runtime libraries from the same build. Nevertheless, in order
 to better support third-party libraries and toolchains that depend on existing
 libomptarget entry points, contributors are discouraged from making
 modifications to them.
+
+Q: Can I use libc functions on the GPU?
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+LLVM provides basic ``libc`` functionality through the LLVM C Library. For 
+building instructions, refer to the associated `LLVM libc documentation 
+<https://libc.llvm.org/gpu/using.html#building-the-gpu-library>`_. Once built, 
+this provides a static library called ``libcgpu.a``. See the documentation for a 
+list of `supported functions <https://libc.llvm.org/gpu/support.html>`_ as well. 
+To utilize these functions, simply link this library as any other when building 
+with OpenMP.
+
+.. code-block:: shell
+
+   clang++ openmp.cpp -fopenmp --offload-arch=gfx90a -lcgpu
+
+For more information on how this is implemented in LLVM/OpenMP's offloading 
+runtime, refer to the `runtime documentation <libomptarget_libc>`_.

openmp/docs/design/Runtimes.rst

@@ -1315,6 +1315,54 @@ LIBOMPTARGET_RPC_LATENCY
 """"""""""""""""""""""""
 This is the maximum amount of time the client will wait for a response from the server.
 
+
+.. _libomptarget_libc:
+
+LLVM/OpenMP support for C library routines
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Support for calling standard C library routines on GPU targets is provided by 
+the `LLVM C Library <https://libc.llvm.org/gpu/>`_. This project provides two 
+static libraries, ``libcgpu.a`` and ``libllvmlibc_rpc_server.a``, which are used 
+by the OpenMP runtime to provide ``libc`` support. The ``libcgpu.a`` library 
+contains the GPU device code, while ``libllvmlibc_rpc_server.a`` provides the 
+interface to the RPC interface. More information on the RPC construction can be 
+found in the `associated documentation <https://libc.llvm.org/gpu/rpc.html>`_.
+
+To provide host services, we run an RPC server inside of the runtime. This 
+allows the host to respond to requests made from the GPU asynchronously. For 
+``libc`` calls that require an RPC server, such as printing, an external handle 
+to the RPC client running on the GPU will be present in the GPU executable. If 
+we find this symbol, we will initialize a client and server and run it in the 
+background while the kernel is executing.
+
+For example, consider the following simple OpenMP offloading code. Here we will 
+simply print a string to the user from the GPU.
+
+.. code-block:: c++
+
+   #include <stdio.h>
+
+   int main() {
+    #pragma omp target
+      { fputs("Hello World!\n", stderr); }
+   }
+
+We can compile this using the ``libcgpu.a`` library to resolve the symbols. 
+Because this function requires RPC support, this will also pull in an externally 
+visible symbol called ``__llvm_libc_rpc_client`` into the device image. When 
+loading the device image, the runtime will check for this symbol and initialize 
+an RPC interface if it is found. The following example shows the RPC server 
+being used.
+
+.. code-block:: console
+
+    $ clang++ hello.c -fopenmp --offload-arch=gfx90a -lcgpu
+    $ env LIBOMPTARGET_DEBUG=1 ./a.out
+    PluginInterface --> Running an RPC server on device 0
+    ...
+    Hello World!
+
 .. _libomptarget_device:
 
 LLVM/OpenMP Target Device Runtime (``libomptarget-ARCH-SUBARCH.bc``)