[Openmp-commits] [PATCH] D155318: [OpenMP] Add documentation on using the `libc` in OpenMP

[Joseph Huber via Phabricator via Openmp-commits]

jhuber6 created this revision.
jhuber6 added reviewers: JonChesterfield, tianshilei1992, jdoerfert.
Herald added subscribers: sunshaoce, guansong, yaxunl.
Herald added a project: All.
jhuber6 requested review of this revision.
Herald added subscribers: openmp-commits, jplehr, sstefan1.
Herald added a project: OpenMP.

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


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D155318

Files:
  openmp/docs/SupportAndFAQ.rst
  openmp/docs/design/Runtimes.rst


Index: openmp/docs/design/Runtimes.rst
===================================================================
--- openmp/docs/design/Runtimes.rst
+++ openmp/docs/design/Runtimes.rst
@@ -1315,6 +1315,58 @@
 """"""""""""""""""""""""
 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. If an RPC server is present, we will 
+run the server instead of blocking during synchronization events. 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!
+
+This interface allows us to provide generic host services to the GPU and can 
+provide the basis of OpenMP reverse offloading.
+
 .. _libomptarget_device:
 
 LLVM/OpenMP Target Device Runtime (``libomptarget-ARCH-SUBARCH.bc``)
Index: openmp/docs/SupportAndFAQ.rst
===================================================================
--- openmp/docs/SupportAndFAQ.rst
+++ openmp/docs/SupportAndFAQ.rst
@@ -433,3 +433,19 @@
 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``. 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>`_.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D155318.540501.patch
Type: text/x-patch
Size: 3757 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20230714/32c2c4d9/attachment.bin>