[OpenMP][NFC] Change error message on offloading failure to mention documentation

[jhuber6]

Summary

This patch changes the error message to instead mention the
documentation page for the debugging options provided by libomptarget
and the bitcode runtimes. Add some extra information to the documentation to
help users more quickly identify debugging resources.

openmp/docs/design/Runtimes.rst

 .. _openmp_runtimes:
 
 LLVM/OpenMP Runtimes
 ====================
 
-There are four distinct types of LLVM/OpenMP runtimes
+There are four distinct types of LLVM/OpenMP runtimes: the host runtime
+:ref:`libomp`, the target offloading runtime :ref:`libomptarget`, the target
+offloading plugin :ref:`libomptarget_plugin`, and finally the target device
+runtime :ref:`libomptarget_device`.
+
+For general information on debugging OpenMP target offloading applications, see
+:ref:`libomptarget_info` and :ref:`libomptarget_device_debugging`
+
+.. _libomp:
 
 LLVM/OpenMP Host Runtime (``libomp``)
 -------------------------------------
 
 An `early (2015) design document <https://openmp.llvm.org/Reference.pdf>`_ for
 the LLVM/OpenMP host runtime, aka.  `libomp.so`, is available as a `pdf
 <https://openmp.llvm.org/Reference.pdf>`_.
 

 KMP_WARNINGS
 """"""""""""
 
 Enables (``true``) or disables (``false``) displaying warnings from the
 OpenMP run-time library during program execution.
 
 **Default:** ``true``
 
+.. _libomptarget:
+
 LLVM/OpenMP Target Host Runtime (``libomptarget``)
 --------------------------------------------------
 
 .. _libopenmptarget_environment_vars:
 
 Environment Variables
 ^^^^^^^^^^^^^^^^^^^^^
 

 
 ``LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD`` sets the threshold size for which the
 ``libomptarget`` memory manager will handle the allocation. Any allocations
 larger than this threshold will not use the memory manager and be freed after
 the device kernel exits. The default threshold value is ``8KB``. If
 ``LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD`` is set to ``0`` the memory manager
 will be completely disabled.
 
+.. _libomptarget_info:
+
 LIBOMPTARGET_INFO
 """""""""""""""""
 
 ``LIBOMPTARGET_INFO`` allows the user to request different types of runtime
 information from ``libomptarget``. ``LIBOMPTARGET_INFO`` uses a 32-bit field to
 enable or disable different types of information. This includes information
 about data-mappings and kernel execution. It is recommended to build your
 application with debugging information enabled, this will enable filenames and
.. code-block:: console
     $ ./sum
 
 .. code-block:: text
 
     CUDA error: an illegal memory access was encountered
     Libomptarget error: Copying data from device failed.
     Libomptarget error: Call to targetDataEnd failed, abort target.
     Libomptarget error: Failed to process data after launching the kernel.
-    Libomptarget error: Run with LIBOMPTARGET_INFO=4 to dump host-target pointer mappings.
+    Libomptarget error: Consult https://openmp.llvm.org/design/Runtimes.html for debugging options.
     sum.cpp:5:1: Libomptarget error 1: failure of target construct while offloading is mandatory
 
 This shows that there is an illegal memory access occuring inside the OpenMP
 target region once execution has moved to the CUDA device, suggesting a
 segmentation fault. This then causes a chain reaction of failures in
 ``libomptarget``. Another message suggests using the ``LIBOMPTARGET_INFO``
 environment variable as described in :ref:`libopenmptarget_environment_vars`. If
 we do this it will print the sate of the host-target pointer mappings at the

 concurrently map the same memory. If the memory is already mapped and the
 map clauses will only modify the reference counter from a non-zero count to
 another non-zero count, concurrent map clauses are supported regardless of
 this option. To disable forced atomic map clauses use "false"/"FALSE" as the
 value of the ``LIBOMPTARGET_MAP_FORCE_ATOMIC`` environment variable.
 The default behavior of LLVM 14 is to force atomic maps clauses, prior versions
 of LLVM did not.
 
+.. _libomptarget_plugin:
 
 LLVM/OpenMP Target Host Runtime Plugins (``libomptarget.rtl.XXXX``)
 -------------------------------------------------------------------
 
 .. _device_runtime:
 
 
 .. _remote_offloading_plugin:

 LIBOMPTARGET_BLOCK_SIZE
 """""""""""""""""""""""
 This is the maximum size of a single message while streaming data transfers between the two endpoints and can be set for both endpoints.
 
 LIBOMPTARGET_RPC_LATENCY
 """"""""""""""""""""""""
 This is the maximum amount of time the client will wait for a response from the server.
 
+.. _libomptarget_device:
+
 LLVM/OpenMP Target Device Runtime (``libomptarget-ARCH-SUBARCH.bc``)
 --------------------------------------------------------------------
 
 The target device runtime is an LLVM bitcode library that implements OpenMP
 runtime functions on the target device. It is linked with the device code's LLVM
 IR during compilation.
 
+.. _libomptarget_device_debugging:
+
 Debugging
 ^^^^^^^^^
 
 The device runtime supports debugging in the runtime itself. This is configured
 at compile-time using the flag ``-fopenmp-target-debug=<N>`` rather than using a
 separate debugging build. If debugging is not enabled, the debugging paths will
 be considered trivially dead and removed by the compiler with zero overhead.
 Debugging is enabled at runtime by running with the environment variable
 ``LIBOMPTARGET_DEVICE_RTL_DEBUG=<N>`` set. The number set is a 32-bit field used
 to selectively enable and disable different features.  Currently, the following
 debugging features are supported.
 
     * Enable debugging assertions in the device. ``0x01``
     * Enable OpenMP runtime function traces in the device. ``0x2``
+    * Enable diagnosing common problems during offloading . ``0x4``
 
 .. code-block:: c++
 
     void copy(double *X, double *Y) {
     #pragma omp target teams distribute parallel for
       for (std::size_t i = 0; i < N; ++i)
         Y[i] = X[i];
     }

openmp/libomptarget/src/omptarget.cpp

FATAL_MESSAGE0(1, "default offloading policy must be switched to "
                       "mandatory or disabled");
     break;
   case tgt_mandatory:
     if (!Success) {
       if (getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE)
         for (auto &Device : PM->Devices)
           dumpTargetPointerMappings(Loc, *Device);
       else
-        FAILURE_MESSAGE("Run with LIBOMPTARGET_INFO=%d to dump host-target "
-                        "pointer mappings.\n",
-                        OMP_INFOTYPE_DUMP_TABLE);
+        FAILURE_MESSAGE("Consult https://openmp.llvm.org/design/Runtimes.html "
+                        "for debugging options.\n");
 
       SourceInfo info(Loc);
       if (info.isAvailible())
         fprintf(stderr, "%s:%d:%d: ", info.getFilename(), info.getLine(),
                 info.getColumn());
       else
         FAILURE_MESSAGE("Source location information not present. Compile with "
                         "-g or -gline-tables-only.\n");