[llvm-branch-commits] [openmp] [OpenMP][docs] Finish MyST migration for OpenMP docs (PR #222194)
Reid Kleckner via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Sep 15 07:47:17 PDT 2026
https://github.com/rnk updated https://github.com/llvm/llvm-project/pull/222194
>From a92353df3ba026a174f478837531668adf1d0760 Mon Sep 17 00:00:00 2001
From: Reid Kleckner <rkleckner at nvidia.com>
Date: Sat, 5 Sep 2026 04:43:01 +0000
Subject: [PATCH 1/3] [OpenMP][docs] Convert reStructuredText files with
rst2myst
---
openmp/docs/CommandLineArgumentReference.md | 377 ++--
openmp/docs/ReleaseNotes.md | 32 +-
openmp/docs/SupportAndFAQ.md | 339 ++--
openmp/docs/design/GPUGenericMode.md | 4 +-
openmp/docs/design/GPUSPMDMode.md | 4 +-
openmp/docs/design/Offloading.md | 16 +-
openmp/docs/design/Overview.md | 23 +-
openmp/docs/design/Runtimes.md | 1783 +++++++++--------
openmp/docs/index.md | 210 +-
openmp/docs/openacc/OpenMPExtensions.md | 175 +-
openmp/docs/openacc/Overview.md | 17 +-
openmp/docs/optimizations/OpenMPOpt.md | 91 +-
.../OpenMPUnawareOptimizations.md | 10 +-
openmp/docs/optimizations/Overview.md | 26 +-
openmp/docs/remarks/OMP100.md | 9 +-
openmp/docs/remarks/OMP101.md | 6 +-
openmp/docs/remarks/OMP102.md | 6 +-
openmp/docs/remarks/OMP110.md | 88 +-
openmp/docs/remarks/OMP111.md | 71 +-
openmp/docs/remarks/OMP112.md | 102 +-
openmp/docs/remarks/OMP113.md | 96 +-
openmp/docs/remarks/OMP120.md | 119 +-
openmp/docs/remarks/OMP121.md | 89 +-
openmp/docs/remarks/OMP130.md | 38 +-
openmp/docs/remarks/OMP131.md | 60 +-
openmp/docs/remarks/OMP132.md | 62 +-
openmp/docs/remarks/OMP133.md | 100 +-
openmp/docs/remarks/OMP140.md | 60 +-
openmp/docs/remarks/OMP150.md | 42 +-
openmp/docs/remarks/OMP160.md | 54 +-
openmp/docs/remarks/OMP170.md | 44 +-
openmp/docs/remarks/OMP180.md | 40 +-
openmp/docs/remarks/OMP190.md | 12 +-
openmp/docs/remarks/OptimizationRemarks.md | 71 +-
34 files changed, 2174 insertions(+), 2102 deletions(-)
diff --git a/openmp/docs/CommandLineArgumentReference.md b/openmp/docs/CommandLineArgumentReference.md
index 06d568263d8681..3b91d4355464d0 100644
--- a/openmp/docs/CommandLineArgumentReference.md
+++ b/openmp/docs/CommandLineArgumentReference.md
@@ -1,186 +1,255 @@
-OpenMP Command-Line Argument Reference
-======================================
-Welcome to the OpenMP in LLVM command line argument reference. The content is
-not a complete list of arguments but includes the essential command-line
-arguments you may need when compiling and linking OpenMP.
-Section :ref:`general_command_line_arguments` lists OpenMP command line options
-for multicore programming while :ref:`offload_command_line_arguments` lists
+# OpenMP Command-Line Argument Reference
+
+Welcome to the OpenMP in LLVM command line argument reference. The content is
+not a complete list of arguments but includes the essential command-line
+arguments you may need when compiling and linking OpenMP.
+Section {ref}`general_command_line_arguments` lists OpenMP command line options
+for multicore programming while {ref}`offload_command_line_arguments` lists
options relevant to OpenMP target offloading.
-.. _general_command_line_arguments:
+(general-command-line-arguments)=
+
+## OpenMP Command-Line Arguments
-OpenMP Command-Line Arguments
------------------------------
+### `-fopenmp`
-``-fopenmp``
-^^^^^^^^^^^^
-Enable the OpenMP compilation toolchain. The compiler will parse OpenMP
+Enable the OpenMP compilation toolchain. The compiler will parse OpenMP
compiler directives and generate parallel code.
-``-fopenmp-extensions``
-^^^^^^^^^^^^^^^^^^^^^^^
-Enable all ``Clang`` extensions for OpenMP directives and clauses. A list of
-current extensions and their implementation status can be found on the
-`support <https://clang.llvm.org/docs/OpenMPSupport.html#openmp-extensions>`_
+### `-fopenmp-extensions`
+
+Enable all `Clang` extensions for OpenMP directives and clauses. A list of
+current extensions and their implementation status can be found on the
+[support](https://clang.llvm.org/docs/OpenMPSupport.html#openmp-extensions)
page.
-``-fopenmp-simd``
-^^^^^^^^^^^^^^^^^
-This option enables OpenMP only for single instruction, multiple data
+### `-fopenmp-simd`
+
+This option enables OpenMP only for single instruction, multiple data
(SIMD) constructs.
-``-static-openmp``
-^^^^^^^^^^^^^^^^^^
+### `-static-openmp`
+
Use the static OpenMP host runtime while linking.
-``-fopenmp-version=<arg>``
-^^^^^^^^^^^^^^^^^^^^^^^^^^
-Set the OpenMP version to a specific version ``<arg>`` of the OpenMP standard.
-For example, you may use ``-fopenmp-version=45`` to select version 4.5 of
-the OpenMP standard. The default value is ``-fopenmp-version=51`` for ``Clang``.
-
-.. _offload_command_line_arguments:
-
-Offloading Specific Command-Line Arguments
-------------------------------------------
-
-.. _fopenmp-targets:
-
-``-fopenmp-targets``
-^^^^^^^^^^^^^^^^^^^^
-| Specify which OpenMP offloading targets should be supported. For example, you
- may specify ``-fopenmp-targets=amdgcn-amd-amdhsa,nvptx64``. This option is
- often optional when :ref:`offload_arch` is provided.
-| It is also possible to offload to CPU architectures, for instance with
- ``-fopenmp-targets=x86_64-pc-linux-gnu``.
-
-.. _offload_arch:
-
-``--offload-arch``
-^^^^^^^^^^^^^^^^^^
-| Specify the device architecture for OpenMP offloading. For instance
- ``--offload-arch=sm_80`` to target an Nvidia Tesla A100,
- ``--offload-arch=gfx90a`` to target an AMD Instinct MI250X, or
- ``--offload-arch=sm_80,gfx90a`` to target both.
-| It is also possible to specify :ref:`fopenmp-targets` without specifying
- ``--offload-arch``. In that case, the executables ``amdgpu-arch`` or
- ``nvptx-arch`` will be executed as part of the compiler driver to
- detect the device architecture automatically.
-| Finally, the device architecture will also be automatically inferred with
- ``--offload-arch=native``.
-
-``--offload-device-only``
-^^^^^^^^^^^^^^^^^^^^^^^^^
-Compile only the code that goes on the device. This option is mainly for
-debugging purposes. It is primarily used for inspecting the intermediate
-representation (IR) output when compiling for the device. It may also be used
+### `-fopenmp-version=<arg>`
+
+Set the OpenMP version to a specific version `<arg>` of the OpenMP standard.
+For example, you may use `-fopenmp-version=45` to select version 4.5 of
+the OpenMP standard. The default value is `-fopenmp-version=51` for `Clang`.
+
+(offload-command-line-arguments)=
+
+## Offloading Specific Command-Line Arguments
+
+(fopenmp-targets)=
+
+### `-fopenmp-targets`
+
+Specify which OpenMP offloading targets should be supported. For example, you
+may specify
+
+`-fopenmp-targets=amdgcn-amd-amdhsa,nvptx64`
+
+. This option is
+often optional when
+
+{ref}`offload_arch`
+
+ is provided.
+
+It is also possible to offload to CPU architectures, for instance with
+
+
+`-fopenmp-targets=x86_64-pc-linux-gnu`
+
+.
+
+(offload-arch)=
+
+### `--offload-arch`
+
+Specify the device architecture for OpenMP offloading. For instance
+
+
+`--offload-arch=sm_80`
+
+ to target an Nvidia Tesla A100,
+
+
+`--offload-arch=gfx90a`
+
+ to target an AMD Instinct MI250X, or
+
+
+`--offload-arch=sm_80,gfx90a`
+
+ to target both.
+
+It is also possible to specify
+
+{ref}`fopenmp-targets`
+
+ without specifying
+
+
+`--offload-arch`
+
+. In that case, the executables
+
+`amdgpu-arch`
+
+ or
+
+
+`nvptx-arch`
+
+ will be executed as part of the compiler driver to
+detect the device architecture automatically.
+
+Finally, the device architecture will also be automatically inferred with
+
+
+`--offload-arch=native`
+
+.
+
+### `--offload-device-only`
+
+Compile only the code that goes on the device. This option is mainly for
+debugging purposes. It is primarily used for inspecting the intermediate
+representation (IR) output when compiling for the device. It may also be used
if device-only runtimes are created.
-``--offload-host-only``
-^^^^^^^^^^^^^^^^^^^^^^^
+### `--offload-host-only`
+
Compile only the code that goes on the host. With this option enabled, the
-``.llvm.offloading`` section with embedded device code will not be included in
+`.llvm.offloading` section with embedded device code will not be included in
the intermediate representation.
-``--offload-host-device``
-^^^^^^^^^^^^^^^^^^^^^^^^^
-Compile the target regions for both the host and the device. That is the
+### `--offload-host-device`
+
+Compile the target regions for both the host and the device. That is the
default option.
-``-Xopenmp-target <arg>``
-^^^^^^^^^^^^^^^^^^^^^^^^^
-Pass an argument ``<arg>`` to the offloading toolchain, for instance
-``-Xopenmp-target -march=sm_80``.
-
-``-Xopenmp-target=<triple> <arg>``
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-Pass an argument ``<arg>`` to the offloading toolchain for the target
-``<triple>``. That is especially useful when an argument must differ for each
-triple. For instance ``-Xopenmp-target=nvptx64 --offload-arch=sm_80
--Xopenmp-target=amdgcn --offload-arch=gfx90a`` to specify the device
-architecture. Alternatively, :ref:`Xarch_host` and :ref:`Xarch_device` can
+### `-Xopenmp-target <arg>`
+
+Pass an argument `<arg>` to the offloading toolchain, for instance
+`-Xopenmp-target -march=sm_80`.
+
+### `-Xopenmp-target=<triple> <arg>`
+
+Pass an argument `<arg>` to the offloading toolchain for the target
+`<triple>`. That is especially useful when an argument must differ for each
+triple. For instance `-Xopenmp-target=nvptx64 --offload-arch=sm_80
+-Xopenmp-target=amdgcn --offload-arch=gfx90a` to specify the device
+architecture. Alternatively, {ref}`Xarch_host` and {ref}`Xarch_device` can
pass an argument to the host and device compilation toolchain.
-``-Xoffload-linker<triple> <arg>``
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-Pass an argument ``<arg>`` to the offloading linker for the target specified in
-``<triple>``.
-
-.. _Xarch_device:
-
-``-Xarch_device <arg>``
-^^^^^^^^^^^^^^^^^^^^^^^
-Pass an argument ``<arg>`` to the device compilation toolchain.
-
-.. _Xarch_host:
-
-``-Xarch_host <arg>``
-^^^^^^^^^^^^^^^^^^^^^
-Pass an argument ``<arg>`` to the host compilation toolchain.
-
-``-foffload-lto[=<arg>]``
-^^^^^^^^^^^^^^^^^^^^^^^^^
-Enable device link time optimization (LTO) and select the LTO mode ``<arg>``.
-Select either ``-foffload-lto=thin`` or ``-foffload-lto=full``. Thin LTO takes
-less time while still achieving some performance gains. If no argument is set,
-this option defaults to ``-foffload-lto=full``.
-
-``-fopenmp-offload-mandatory``
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-| This option is set to avoid generating the host fallback code
- executed when offloading to the device fails. That is
- helpful when the target contains code that cannot be compiled for the host, for
- instance, if it contains unguarded device intrinsics.
-| This option can also be used to reduce compile time.
-| This option should not be used when one wants to verify that the code is being
- offloaded to the device. Instead, set the environment variable
- ``OMP_TARGET_OFFLOAD='MANDATORY'`` to confirm that the code is being offloaded to
- the device.
-
-``-fopenmp-target-debug[=<arg>]``
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-Enable debugging in the device runtime library (RTL). Note that it is both
-necessary to configure the debugging in the device runtime at compile-time with
-``-fopenmp-target-debug=<arg>`` and enable debugging at runtime with the
-environment variable ``LIBOMPTARGET_DEVICE_RTL_DEBUG=<arg>``. Further, it is
-currently only supported for Nvidia targets as of July 2023. Alternatively, the
-environment variable ``LIBOMPTARGET_DEBUG`` can be set to debug both Nvidia and
-AMD GPU targets. For more information, see the
-`debugging instructions <https://openmp.llvm.org/design/Runtimes.html#debugging>`_.
+### `-Xoffload-linker<triple> <arg>`
+
+Pass an argument `<arg>` to the offloading linker for the target specified in
+`<triple>`.
+
+(xarch-device)=
+
+### `-Xarch_device <arg>`
+
+Pass an argument `<arg>` to the device compilation toolchain.
+
+(xarch-host)=
+
+### `-Xarch_host <arg>`
+
+Pass an argument `<arg>` to the host compilation toolchain.
+
+### `-foffload-lto[=<arg>]`
+
+Enable device link time optimization (LTO) and select the LTO mode `<arg>`.
+Select either `-foffload-lto=thin` or `-foffload-lto=full`. Thin LTO takes
+less time while still achieving some performance gains. If no argument is set,
+this option defaults to `-foffload-lto=full`.
+
+### `-fopenmp-offload-mandatory`
+
+This option is set to avoid generating the host fallback code
+executed when offloading to the device fails. That is
+helpful when the target contains code that cannot be compiled for the host, for
+instance, if it contains unguarded device intrinsics.
+
+This option can also be used to reduce compile time.
+
+This option should not be used when one wants to verify that the code is being
+offloaded to the device. Instead, set the environment variable
+
+
+`OMP_TARGET_OFFLOAD='MANDATORY'`
+
+ to confirm that the code is being offloaded to
+the device.
+
+### `-fopenmp-target-debug[=<arg>]`
+
+Enable debugging in the device runtime library (RTL). Note that it is both
+necessary to configure the debugging in the device runtime at compile-time with
+`-fopenmp-target-debug=<arg>` and enable debugging at runtime with the
+environment variable `LIBOMPTARGET_DEVICE_RTL_DEBUG=<arg>`. Further, it is
+currently only supported for Nvidia targets as of July 2023. Alternatively, the
+environment variable `LIBOMPTARGET_DEBUG` can be set to debug both Nvidia and
+AMD GPU targets. For more information, see the
+[debugging instructions](https://openmp.llvm.org/design/Runtimes.html#debugging).
The debugging instructions list the supported debugging arguments.
-``-fopenmp-target-jit``
-^^^^^^^^^^^^^^^^^^^^^^^
-| Emit code that is Just-in-Time (JIT) compiled for OpenMP offloading. Embed
- LLVM-IR for the device code in the object files rather than binary code for the
- respective target. At runtime, the LLVM-IR is optimized again and compiled for
- the target device. The optimization level can be set at runtime with
- ``LIBOMPTARGET_JIT_OPT_LEVEL``, for instance,
- ``LIBOMPTARGET_JIT_OPT_LEVEL=3`` corresponding to optimizations level ``-O3``.
- See the
- `OpenMP JIT details <https://openmp.llvm.org/design/Runtimes.html#libomptarget-jit-pre-opt-ir-module>`_
- for instructions on extracting the embedded device code before or after the
- JIT and more.
-| We want to emphasize that JIT for OpenMP offloading is good for debugging as
- the target IR can be extracted, modified, and injected at runtime.
-
-``--offload-new-driver``
-^^^^^^^^^^^^^^^^^^^^^^^^
+### `-fopenmp-target-jit`
+
+Emit code that is Just-in-Time (JIT) compiled for OpenMP offloading. Embed
+LLVM-IR for the device code in the object files rather than binary code for the
+respective target. At runtime, the LLVM-IR is optimized again and compiled for
+the target device. The optimization level can be set at runtime with
+
+
+`LIBOMPTARGET_JIT_OPT_LEVEL`
+
+, for instance,
+
+
+`LIBOMPTARGET_JIT_OPT_LEVEL=3`
+
+ corresponding to optimizations level
+
+`-O3`
+
+.
+See the
+
+
+[OpenMP JIT details](https://openmp.llvm.org/design/Runtimes.html#libomptarget-jit-pre-opt-ir-module)
+
+
+for instructions on extracting the embedded device code before or after the
+JIT and more.
+
+We want to emphasize that JIT for OpenMP offloading is good for debugging as
+the target IR can be extracted, modified, and injected at runtime.
+
+### `--offload-new-driver`
+
Deprecated and accepted as a no-op. The new offloading driver is always used for
offloading compilation; the legacy driver has been removed.
-``--offload-link``
-^^^^^^^^^^^^^^^^^^
-Use the new offloading linker `clang-linker-wrapper` to perform the link job.
-`clang-linker-wrapper` is the default offloading linker for OpenMP. This option
-can be used to use the new offloading linker in toolchains that do not automatically
+### `--offload-link`
+
+Use the new offloading linker `clang-linker-wrapper` to perform the link job.
+`clang-linker-wrapper` is the default offloading linker for OpenMP. This option
+can be used to use the new offloading linker in toolchains that do not automatically
use it. It is necessary to enable this option when linking with CUDA or HIP files.
-``-nogpulib``
-^^^^^^^^^^^^^
+### `-nogpulib`
+
Do not link the device library for CUDA or HIP device compilation.
-``-nogpuinc``
-^^^^^^^^^^^^^
+### `-nogpuinc`
+
Do not include the default CUDA or HIP headers, and do not add CUDA or HIP
include paths.
+
diff --git a/openmp/docs/ReleaseNotes.md b/openmp/docs/ReleaseNotes.md
index fcf216d1838006..ba597999938acd 100644
--- a/openmp/docs/ReleaseNotes.md
+++ b/openmp/docs/ReleaseNotes.md
@@ -1,33 +1,29 @@
-===========================
-OpenMP 21.0.0 Release Notes
-===========================
+# OpenMP 21.0.0 Release Notes
+:::{warning}
+These are in-progress notes for the upcoming LLVM 19.0.0 release.
+Release notes for previous releases can be found on
+[the Download Page](https://releases.llvm.org/download.html).
+:::
-.. warning::
- These are in-progress notes for the upcoming LLVM 19.0.0 release.
- Release notes for previous releases can be found on
- `the Download Page <https://releases.llvm.org/download.html>`_.
-
-
-Introduction
-============
+## Introduction
This document contains the release notes for the OpenMP runtime, release 19.0.0.
Here we describe the status of OpenMP, including major improvements
from the previous release. All OpenMP releases may be downloaded
-from the `LLVM releases web site <https://llvm.org/releases/>`_.
+from the [LLVM releases web site](https://llvm.org/releases/).
-Non-comprehensive list of changes in this release
-=================================================
+## Non-comprehensive list of changes in this release
- Removed the standalone build mode. It is redundant with the runtimes default
build.
-Device Runtime
---------------
+### Device Runtime
+
- Changed the OpenMP DeviceRTL to use 'generic' IR. The
- ``LIBOMPTARGET_DEVICE_ARCHITECTURES`` CMake argument is now unused and will
+ `LIBOMPTARGET_DEVICE_ARCHITECTURES` CMake argument is now unused and will
always build support for AMDGPU and NVPTX targets.
- Updated the offloading entry format but retained backwards compatibility with
the old format.
-- The LLVM_ENABLE_PROJECTS=openmp build mode has been removed.
\ No newline at end of file
+- The LLVM_ENABLE_PROJECTS=openmp build mode has been removed.
+
diff --git a/openmp/docs/SupportAndFAQ.md b/openmp/docs/SupportAndFAQ.md
index 3ee7a75dc3d2ef..ee25f9566faee4 100644
--- a/openmp/docs/SupportAndFAQ.md
+++ b/openmp/docs/SupportAndFAQ.md
@@ -1,107 +1,97 @@
-Support, Getting Involved, and FAQ
-==================================
+# Support, Getting Involved, and FAQ
-Please do not hesitate to reach out to us on the `Discourse forums (Runtimes - OpenMP) <https://discourse.llvm.org/c/runtimes/openmp/35>`_ or join
-one of our :ref:`regular calls <calls>`. Some common questions are answered in
-the :ref:`faq`.
+Please do not hesitate to reach out to us on the [Discourse forums (Runtimes - OpenMP)](https://discourse.llvm.org/c/runtimes/openmp/35) or join
+one of our {ref}`regular calls <calls>`. Some common questions are answered in
+the {ref}`faq`.
-.. _calls:
+(calls)=
-Calls
------
+## Calls
-OpenMP in LLVM Technical Call
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### OpenMP in LLVM Technical Call
-- Development updates on OpenMP (and OpenACC) in the LLVM Project, including Clang, optimization, and runtime work.
-- Join `OpenMP in LLVM Technical Call <https://bluejeans.com/544112769//webrtc>`__.
-- Time: Weekly call on every Wednesday 7:00 AM Pacific time.
-- Meeting minutes are `here <https://docs.google.com/document/d/1Tz8WFN13n7yJ-SCE0Qjqf9LmjGUw0dWO9Ts1ss4YOdg/edit>`__.
-- Status tracking `page <https://openmp.llvm.org/docs>`__.
+- Development updates on OpenMP (and OpenACC) in the LLVM Project, including Clang, optimization, and runtime work.
+- Join [OpenMP in LLVM Technical Call](https://bluejeans.com/544112769//webrtc).
+- Time: Weekly call on every Wednesday 7:00 AM Pacific time.
+- Meeting minutes are [here](https://docs.google.com/document/d/1Tz8WFN13n7yJ-SCE0Qjqf9LmjGUw0dWO9Ts1ss4YOdg/edit).
+- Status tracking [page](https://openmp.llvm.org/docs).
+### OpenMP in Flang Technical Call
-OpenMP in Flang Technical Call
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-- Development updates on OpenMP and OpenACC in the Flang Project.
-- Join `OpenMP in Flang Technical Call <https://bit.ly/39eQW3o>`_
-- Time: Weekly call on every Thursdays 8:00 AM Pacific time.
-- Meeting minutes are `here <https://docs.google.com/document/d/1yA-MeJf6RYY-ZXpdol0t7YoDoqtwAyBhFLr5thu5pFI>`__.
-- Status tracking `page <https://flang.llvm.org/docs/OpenMPSupport.html>`__.
+- Development updates on OpenMP and OpenACC in the Flang Project.
+- Join [OpenMP in Flang Technical Call](https://bit.ly/39eQW3o)
+- Time: Weekly call on every Thursdays 8:00 AM Pacific time.
+- Meeting minutes are [here](https://docs.google.com/document/d/1yA-MeJf6RYY-ZXpdol0t7YoDoqtwAyBhFLr5thu5pFI).
+- Status tracking [page](https://flang.llvm.org/docs/OpenMPSupport.html).
+(faq)=
-.. _faq:
+## FAQ
-FAQ
----
+:::{note}
+The FAQ is a work in progress and most of the expected content is not
+yet available. While you can expect changes, we always welcome feedback and
+additions. Please post on the [Discourse forums (Runtimes - OpenMP)](https://discourse.llvm.org/c/runtimes/openmp/35).
+:::
-.. note::
- The FAQ is a work in progress and most of the expected content is not
- yet available. While you can expect changes, we always welcome feedback and
- additions. Please post on the `Discourse forums (Runtimes - OpenMP) <https://discourse.llvm.org/c/runtimes/openmp/35>`__.
+### Q: How to contribute a patch to the webpage or any other part?
+All patches go through the regular [LLVM review process](https://llvm.org/docs/Contributing.html#how-to-submit-a-patch).
-Q: How to contribute a patch to the webpage or any other part?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-
-All patches go through the regular `LLVM review process
-<https://llvm.org/docs/Contributing.html#how-to-submit-a-patch>`_.
-
-
-Q: What are the known limitations of OpenMP AMDGPU offload?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: What are the known limitations of OpenMP AMDGPU offload?
LD_LIBRARY_PATH or rpath/runpath are required to find libomp.so and
libomptarget.so correctly. The recommended way to configure this is with the
-``-frtlib-add-rpath`` option. Alternatively, set the ``LD_LIBRARY_PATH``
+`-frtlib-add-rpath` option. Alternatively, set the `LD_LIBRARY_PATH`
environment variable to point to the installation. Normally, these libraries are
installed in the target specific runtime directory. For example, a typical
installation will have
-``<install>/lib/x86_64-unknown-linux-gnu/llibomptarget.so``
+`<install>/lib/x86_64-unknown-linux-gnu/llibomptarget.so`
Some versions of the driver for the radeon vii (gfx906) will error unless the
environment variable 'export HSA_IGNORE_SRAMECC_MISREPORT=1' is set.
-Q: What are the LLVM components used in offloading and how are they found?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: What are the LLVM components used in offloading and how are they found?
+
The libraries used by an executable compiled for target offloading are:
-- ``libomp.so`` (or similar), the host openmp runtime
-- ``libomptarget.so``, the target-agnostic target offloading openmp runtime
-- ``libompdevice.a``, the device-side OpenMP runtime.
+- `libomp.so` (or similar), the host openmp runtime
+- `libomptarget.so`, the target-agnostic target offloading openmp runtime
+- `libompdevice.a`, the device-side OpenMP runtime.
- dependencies of those plugins, e.g. cuda/rocr for nvptx/amdgpu
The compiled executable is dynamically linked against a host runtime, e.g.
-``libomp.so``, and against the target offloading runtime, ``libomptarget.so``. These
+`libomp.so`, and against the target offloading runtime, `libomptarget.so`. These
are found like any other dynamic library, by setting rpath or runpath on the
-executable, by setting ``LD_LIBRARY_PATH``, or by adding them to the system search.
-
-``libomptarget.so`` is only supported to work with the associated ``clang``
-compiler. On systems with globally installed ``libomptarget.so`` this can be
-problematic. For this reason it is recommended to use a `Clang configuration
-file <https://clang.llvm.org/docs/UsersManual.html#configuration-files>`__ to
-automatically configure the environment. For example, store the following file
-as ``openmp.cfg`` next to your ``clang`` executable.
+executable, by setting `LD_LIBRARY_PATH`, or by adding them to the system search.
-.. code-block:: text
+`libomptarget.so` is only supported to work with the associated `clang`
+compiler. On systems with globally installed `libomptarget.so` this can be
+problematic. For this reason it is recommended to use a [Clang configuration
+file](https://clang.llvm.org/docs/UsersManual.html#configuration-files) to
+automatically configure the environment. For example, store the following file
+as `openmp.cfg` next to your `clang` executable.
- # Library paths for OpenMP offloading.
- -L '<CFGDIR>/../lib'
- -Wl,-rpath='<CFGDIR>/../lib'
+```text
+# Library paths for OpenMP offloading.
+-L '<CFGDIR>/../lib'
+-Wl,-rpath='<CFGDIR>/../lib'
+```
The plugins will try to find their dependencies in plugin-dependent fashion.
The cuda plugin is dynamically linked against libcuda if cmake found it at
-compiler build time. Otherwise it will attempt to dlopen ``libcuda.so``. It does
+compiler build time. Otherwise it will attempt to dlopen `libcuda.so`. It does
not have rpath set.
The amdgpu plugin is linked against ROCr if cmake found it at compiler build
-time. Otherwise it will attempt to dlopen ``libhsa-runtime64.so``. It has rpath
-set to ``$ORIGIN``, so installing ``libhsa-runtime64.so`` in the same directory is a
+time. Otherwise it will attempt to dlopen `libhsa-runtime64.so`. It has rpath
+set to `$ORIGIN`, so installing `libhsa-runtime64.so` in the same directory is a
way to locate it without environment variables.
In addition to those, there is a compiler runtime library called deviceRTL.
This is compiled from mostly common code into an architecture specific
-bitcode library, e.g. ``libomptarget-nvptx-sm_70.bc``.
+bitcode library, e.g. `libomptarget-nvptx-sm_70.bc`.
Clang and the deviceRTL need to match closely as the interface between them
changes frequently. Using both from the same monorepo checkout is strongly
@@ -109,24 +99,22 @@ recommended.
Unlike the host side which lets environment variables select components, the
deviceRTL that is located in the clang lib directory is preferred. Only if
-it is absent, the ``LIBRARY_PATH`` environment variable is searched to find a
+it is absent, the `LIBRARY_PATH` environment variable is searched to find a
bitcode file with the right name. This can be overridden by passing a clang
-flag, ``--libomptarget-nvptx-bc-path`` or ``--libomptarget-amdgcn-bc-path``. That
+flag, `--libomptarget-nvptx-bc-path` or `--libomptarget-amdgcn-bc-path`. That
can specify a directory or an exact bitcode file to use.
+### Q: Does OpenMP offloading support work in pre-packaged LLVM releases?
+
+For now, the answer is most likely *no*. Please see {ref}`build_offload_capable_compiler`.
-Q: Does OpenMP offloading support work in pre-packaged LLVM releases?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-For now, the answer is most likely *no*. Please see :ref:`build_offload_capable_compiler`.
+### Q: Does OpenMP offloading support work in packages distributed as part of my OS?
-Q: Does OpenMP offloading support work in packages distributed as part of my OS?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-For now, the answer is most likely *no*. Please see :ref:`build_offload_capable_compiler`.
+For now, the answer is most likely *no*. Please see {ref}`build_offload_capable_compiler`.
-.. _math_and_complex_in_target_regions:
+(math-and-complex-in-target-regions)=
-Q: Does Clang support `<math.h>` and `<complex.h>` operations in OpenMP target on GPUs?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: Does Clang support `<math.h>` and `<complex.h>` operations in OpenMP target on GPUs?
Yes, LLVM/Clang allows math functions and complex arithmetic inside of OpenMP
target regions that are compiled for GPUs.
@@ -147,34 +135,29 @@ vendor, in an OpenMP begin/end declare variant. These functions will then be
picked up instead of the host versions while host only variables and function
definitions are still available. Complex arithmetic and functions are support
through a similar mechanism. It is worth noting that this support requires
-`extensions to the OpenMP begin/end declare variant context selector
-<https://clang.llvm.org/docs/AttributeReference.html#pragma-omp-declare-variant>`__
+[extensions to the OpenMP begin/end declare variant context selector](https://clang.llvm.org/docs/AttributeReference.html#pragma-omp-declare-variant)
that are exposed through LLVM/Clang to the user as well.
-Q: Can I use dynamically linked libraries with OpenMP offloading?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: Can I use dynamically linked libraries with OpenMP offloading?
Dynamically linked libraries can be used if there is no device code shared
between the library and application. Anything declared on the device inside the
shared library will not be visible to the application when it's linked. This is
because device code only supports static linking.
-Q: How to build an OpenMP offload capable compiler with an outdated host compiler?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: How to build an OpenMP offload capable compiler with an outdated host compiler?
Enabling the OpenMP runtime will perform a two-stage build for you.
If your host compiler is different from your system-wide compiler, you may need
-to set ``CMAKE_{C,CXX}_FLAGS`` like
-``--gcc-install-dir=/usr/lib/gcc/x86_64-linux-gnu/12`` so that clang will be
+to set `CMAKE_{C,CXX}_FLAGS` like
+`--gcc-install-dir=/usr/lib/gcc/x86_64-linux-gnu/12` so that clang will be
able to find the correct GCC toolchain in the second stage of the build.
For example, if your system-wide GCC installation is too old to build LLVM and
-you would like to use a newer GCC, set ``--gcc-install-dir=``
+you would like to use a newer GCC, set `--gcc-install-dir=`
to inform clang of the GCC installation you would like to use in the second stage.
-
-Q: What does 'Stack size for entry function cannot be statically determined' mean?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: What does 'Stack size for entry function cannot be statically determined' mean?
This is a warning that the Nvidia tools will sometimes emit if the offloading
region is too complex. Normally, the CUDA tools attempt to statically determine
@@ -183,11 +166,10 @@ thread will have as much memory as it needs. If the control flow of the kernel
is too complex, containing recursive calls or nested parallelism, this analysis
can fail. If this warning is triggered it means that the kernel may run out of
stack memory during execution and crash. The environment variable
-``LIBOMPTARGET_STACK_SIZE`` can be used to increase the stack size if this
+`LIBOMPTARGET_STACK_SIZE` can be used to increase the stack size if this
occurs.
-Q: Can OpenMP offloading compile for multiple architectures?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: Can OpenMP offloading compile for multiple architectures?
Since LLVM version 15.0, OpenMP offloading supports offloading to multiple
architectures at once. This allows for executables to be run on different
@@ -196,62 +178,61 @@ multiple sub-architectures for the same target. Additionally, static libraries
will only extract archive members if an architecture is used, allowing users to
create generic libraries.
-The architecture can either be specified manually using ``--offload-arch=``. If
-``--offload-arch=`` is present and no ``-fopenmp-targets=`` flag is present then
+The architecture can either be specified manually using `--offload-arch=`. If
+`--offload-arch=` is present and no `-fopenmp-targets=` flag is present then
the targets will be inferred from the architectures. Conversely, if
-``--fopenmp-targets=`` is present with no ``--offload-arch`` then the target
+`--fopenmp-targets=` is present with no `--offload-arch` then the target
architecture will be set to a default value, usually the architecture supported
-by the system LLVM was built on by executing the ``offload-arch`` utility.
+by the system LLVM was built on by executing the `offload-arch` utility.
For example, an executable can be built that runs on AMDGPU and NVIDIA hardware
given that the necessary build tools are installed for both.
-.. code-block:: shell
-
- clang example.c -fopenmp --offload-arch=gfx90a --offload-arch=sm_80
+```shell
+clang example.c -fopenmp --offload-arch=gfx90a --offload-arch=sm_80
+```
If just given the architectures we should be able to infer the triples,
otherwise we can specify them manually.
-.. code-block:: shell
-
- clang example.c -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda \
- -Xopenmp-target=amdgcn-amd-amdhsa --offload-arch=gfx90a \
- -Xopenmp-target=nvptx64-nvidia-cuda --offload-arch=sm_80
+```shell
+clang example.c -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda \
+ -Xopenmp-target=amdgcn-amd-amdhsa --offload-arch=gfx90a \
+ -Xopenmp-target=nvptx64-nvidia-cuda --offload-arch=sm_80
+```
When linking against a static library that contains device code for multiple
architectures, only the images used by the executable will be extracted.
-.. code-block:: shell
-
- clang example.c -fopenmp --offload-arch=gfx90a,gfx90a,sm_70,sm_80 -c
- llvm-ar rcs libexample.a example.o
- clang app.c -fopenmp --offload-arch=gfx90a -o app
-
-The supported device images can be viewed using the ``--offloading`` option with
-``llvm-objdump``.
+```shell
+clang example.c -fopenmp --offload-arch=gfx90a,gfx90a,sm_70,sm_80 -c
+llvm-ar rcs libexample.a example.o
+clang app.c -fopenmp --offload-arch=gfx90a -o app
+```
-.. code-block:: shell
+The supported device images can be viewed using the `--offloading` option with
+`llvm-objdump`.
- clang example.c -fopenmp --offload-arch=gfx90a --offload-arch=sm_80 -o example
- llvm-objdump --offloading example
+```shell
+clang example.c -fopenmp --offload-arch=gfx90a --offload-arch=sm_80 -o example
+llvm-objdump --offloading example
- a.out: file format elf64-x86-64
+a.out: file format elf64-x86-64
- OFFLOADING IMAGE [0]:
- kind elf
- arch gfx90a
- triple amdgcn-amd-amdhsa
- producer openmp
+OFFLOADING IMAGE [0]:
+kind elf
+arch gfx90a
+triple amdgcn-amd-amdhsa
+producer openmp
- OFFLOADING IMAGE [1]:
- kind elf
- arch sm_80
- triple nvptx64-nvidia-cuda
- producer openmp
+OFFLOADING IMAGE [1]:
+kind elf
+arch sm_80
+triple nvptx64-nvidia-cuda
+producer openmp
+```
-Q: Can I link OpenMP offloading with CUDA or HIP?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: Can I link OpenMP offloading with CUDA or HIP?
OpenMP offloading files can currently be experimentally linked with CUDA and HIP
files. This will allow OpenMP to call a CUDA device function or vice-versa.
@@ -259,74 +240,72 @@ However, the global state will be distinct between the two images at runtime.
This means any global variables will potentially have different values when
queried from OpenMP or CUDA.
-Linking CUDA and HIP currently requires linking using ``--offload-link``.
-Additionally, ``-fgpu-rdc`` must be used to create a linkable device image.
+Linking CUDA and HIP currently requires linking using `--offload-link`.
+Additionally, `-fgpu-rdc` must be used to create a linkable device image.
-.. code-block:: shell
+```shell
+clang++ openmp.cpp -fopenmp --offload-arch=sm_80 -c
+clang++ cuda.cu --offload-arch=sm_80 -fgpu-rdc -c
+clang++ openmp.o cuda.o --offload-link -o app
+```
- clang++ openmp.cpp -fopenmp --offload-arch=sm_80 -c
- clang++ cuda.cu --offload-arch=sm_80 -fgpu-rdc -c
- clang++ openmp.o cuda.o --offload-link -o app
-
-Q: Are libomptarget and plugins backward compatible?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### Q: Are libomptarget and plugins backward compatible?
No. libomptarget and plugins are now built as LLVM libraries starting from LLVM
-15. Because LLVM libraries are not backward compatible, libomptarget and plugins
+15\. Because LLVM libraries are not backward compatible, libomptarget and plugins
are not as well. Given that fact, the interfaces between 1) the Clang compiler
and libomptarget, 2) the Clang compiler and device runtime library, and
-3) libomptarget and plugins are not guaranteed to be compatible with an earlier
+3\) libomptarget and plugins are not guaranteed to be compatible with an earlier
version. Users are responsible for ensuring compatibility when not using the
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?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### 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
+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 -Xoffload-linker -lc
-
-For more information on how this is implemented in LLVM/OpenMP's offloading
-runtime, refer to the `runtime documentation <libomptarget_libc>`_.
-
-Q: What command line options can I use for OpenMP?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-We recommend taking a look at the OpenMP
-:doc:`command line argument reference <CommandLineArgumentReference>` page.
-
-Q: Can I build the offloading runtimes without CUDA or HSA?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-By default, the offloading runtime will load the associated vendor runtime
-during initialization rather than directly linking against them. This allows the
-program to be built and run on many machine. If you wish to directly link
-against these libraries, use the ``LIBOMPTARGET_DLOPEN_PLUGINS=""`` option to
-suppress it for each plugin. The default value is every plugin enabled with
-``LIBOMPTARGET_PLUGINS_TO_BUILD``.
-
-Q: Why is my build taking a long time?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-When installing OpenMP and other LLVM components, the build time on multicore
-systems can be significantly reduced with parallel build jobs. As suggested in
-*LLVM Techniques, Tips, and Best Practices*, one could consider using ``ninja`` as the
-generator. This can be done with the CMake option ``cmake -G Ninja``. Afterward,
-use ``ninja install`` and specify the number of parallel jobs with ``-j``. The build
-time can also be reduced by setting the build type to ``Release`` with the
-``CMAKE_BUILD_TYPE`` option. Recompilation can also be sped up by caching previous
-compilations. Consider enabling ``Ccache`` with
-``CMAKE_CXX_COMPILER_LAUNCHER=ccache``.
-
-Q: Did this FAQ not answer your question?
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-Feel free to post questions or browse old threads at
-`LLVM Discourse <https://discourse.llvm.org/c/runtimes/openmp/>`__.
+```shell
+clang++ openmp.cpp -fopenmp --offload-arch=gfx90a -Xoffload-linker -lc
+```
+
+For more information on how this is implemented in LLVM/OpenMP's offloading
+runtime, refer to the [runtime documentation](libomptarget_libc).
+
+### Q: What command line options can I use for OpenMP?
+
+We recommend taking a look at the OpenMP
+{doc}`command line argument reference <CommandLineArgumentReference>` page.
+
+### Q: Can I build the offloading runtimes without CUDA or HSA?
+
+By default, the offloading runtime will load the associated vendor runtime
+during initialization rather than directly linking against them. This allows the
+program to be built and run on many machine. If you wish to directly link
+against these libraries, use the `LIBOMPTARGET_DLOPEN_PLUGINS=""` option to
+suppress it for each plugin. The default value is every plugin enabled with
+`LIBOMPTARGET_PLUGINS_TO_BUILD`.
+
+### Q: Why is my build taking a long time?
+
+When installing OpenMP and other LLVM components, the build time on multicore
+systems can be significantly reduced with parallel build jobs. As suggested in
+*LLVM Techniques, Tips, and Best Practices*, one could consider using `ninja` as the
+generator. This can be done with the CMake option `cmake -G Ninja`. Afterward,
+use `ninja install` and specify the number of parallel jobs with `-j`. The build
+time can also be reduced by setting the build type to `Release` with the
+`CMAKE_BUILD_TYPE` option. Recompilation can also be sped up by caching previous
+compilations. Consider enabling `Ccache` with
+`CMAKE_CXX_COMPILER_LAUNCHER=ccache`.
+
+### Q: Did this FAQ not answer your question?
+
+Feel free to post questions or browse old threads at
+[LLVM Discourse](https://discourse.llvm.org/c/runtimes/openmp/).
+
diff --git a/openmp/docs/design/GPUGenericMode.md b/openmp/docs/design/GPUGenericMode.md
index 719e2d50202552..143e657ace508c 100644
--- a/openmp/docs/design/GPUGenericMode.md
+++ b/openmp/docs/design/GPUGenericMode.md
@@ -1,2 +1,2 @@
-OpenMP Target Offloading --- Generic Mode
-=========================================
+# OpenMP Target Offloading --- Generic Mode
+
diff --git a/openmp/docs/design/GPUSPMDMode.md b/openmp/docs/design/GPUSPMDMode.md
index ea55ffa6bc8ea7..6f2b30317ae88a 100644
--- a/openmp/docs/design/GPUSPMDMode.md
+++ b/openmp/docs/design/GPUSPMDMode.md
@@ -1,2 +1,2 @@
-OpenMP Target Offloading --- SPMD Mode
-======================================
+# OpenMP Target Offloading --- SPMD Mode
+
diff --git a/openmp/docs/design/Offloading.md b/openmp/docs/design/Offloading.md
index 3296d0729ffcc8..85af57e4ebf68f 100644
--- a/openmp/docs/design/Offloading.md
+++ b/openmp/docs/design/Offloading.md
@@ -1,11 +1,11 @@
-OpenMP in LLVM --- Offloading Design
-====================================
+# OpenMP in LLVM --- Offloading Design
-.. toctree::
- :glob:
- :hidden:
- :maxdepth: 1
+```{toctree}
+:glob: true
+:hidden: true
+:maxdepth: 1
- GPUSPMDMode
- GPUGenericMode
+GPUSPMDMode
+GPUGenericMode
+```
diff --git a/openmp/docs/design/Overview.md b/openmp/docs/design/Overview.md
index 06c0ac930c30a9..90b0029fda55e3 100644
--- a/openmp/docs/design/Overview.md
+++ b/openmp/docs/design/Overview.md
@@ -1,18 +1,15 @@
-OpenMP in LLVM --- Design Overview
-==================================
+# OpenMP in LLVM --- Design Overview
+## Resources
-Resources
----------
+- OpenMP Booth @ SC19: "OpenMP clang and flang Development" <https://youtu.be/6yOa-hRi63M>
-- OpenMP Booth @ SC19: "OpenMP clang and flang Development" https://youtu.be/6yOa-hRi63M
+```{toctree}
+:glob: true
+:hidden: true
+:maxdepth: 1
-
-.. toctree::
- :glob:
- :hidden:
- :maxdepth: 1
-
- Runtimes
- Offloading
+Runtimes
+Offloading
+```
diff --git a/openmp/docs/design/Runtimes.md b/openmp/docs/design/Runtimes.md
index 0e75925c787426..e31dcaf9d5e9cd 100644
--- a/openmp/docs/design/Runtimes.md
+++ b/openmp/docs/design/Runtimes.md
@@ -1,100 +1,101 @@
-.. _openmp_runtimes:
+(openmp-runtimes)=
-LLVM/OpenMP Runtimes
-====================
+# 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`.
+{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`
+{ref}`libomptarget_info` and {ref}`libomptarget_device_debugging`
-.. _libomp:
+(libomp)=
-LLVM/OpenMP Host Runtime (``libomp``)
--------------------------------------
+## LLVM/OpenMP Host Runtime (`libomp`)
-An `early (2015) design document
-<https://raw.githubusercontent.com/llvm/llvm-project/main/openmp/runtime/doc/Reference.pdf>`_
-for the LLVM/OpenMP host runtime, aka. `libomp.so`, is available as a `pdf
-<https://raw.githubusercontent.com/llvm/llvm-project/main/openmp/runtime/doc/Reference.pdf>`_.
+An [early (2015) design document](https://raw.githubusercontent.com/llvm/llvm-project/main/openmp/runtime/doc/Reference.pdf)
+for the LLVM/OpenMP host runtime, aka. `libomp.so`, is available as a [pdf](https://raw.githubusercontent.com/llvm/llvm-project/main/openmp/runtime/doc/Reference.pdf).
-.. _libomp_environment_vars:
+(libomp-environment-vars)=
-Environment Variables
-^^^^^^^^^^^^^^^^^^^^^
+### Environment Variables
-OMP_CANCELLATION
-""""""""""""""""
+#### OMP_CANCELLATION
Enables cancellation of the innermost enclosing region of the type specified.
-If set to ``true``, the effects of the cancel construct and of cancellation
-points are enabled and cancellation is activated. If set to ``false``,
+If set to `true`, the effects of the cancel construct and of cancellation
+points are enabled and cancellation is activated. If set to `false`,
cancellation is disabled and the cancel construct and cancellation points are
effectively ignored.
-.. note::
- Internal barrier code will work differently depending on whether cancellation
- is enabled. Barrier code should repeatedly check the global flag to figure
- out if cancellation has been triggered. If a thread observes cancellation, it
- should leave the barrier prematurely with the return value 1 (and may wake up
- other threads). Otherwise, it should leave the barrier with the return value 0.
+:::{note}
+Internal barrier code will work differently depending on whether cancellation
+is enabled. Barrier code should repeatedly check the global flag to figure
+out if cancellation has been triggered. If a thread observes cancellation, it
+should leave the barrier prematurely with the return value 1 (and may wake up
+other threads). Otherwise, it should leave the barrier with the return value 0.
+:::
-Enables (``true``) or disables (``false``) cancellation of the innermost
+Enables (`true`) or disables (`false`) cancellation of the innermost
enclosing region of the type specified.
-**Default:** ``false``
+**Default:** `false`
+#### OMP_DISPLAY_ENV
-OMP_DISPLAY_ENV
-"""""""""""""""
-
-Enables (``true``) or disables (``false``) the printing to ``stderr`` of
+Enables (`true`) or disables (`false`) the printing to `stderr` of
the OpenMP version number and the values associated with the OpenMP
environment variables.
-Possible values are: ``true``, ``false``, or ``verbose``.
+Possible values are: `true`, `false`, or `verbose`.
-**Default:** ``false``
+**Default:** `false`
-OMP_DEFAULT_DEVICE
-""""""""""""""""""
+#### OMP_DEFAULT_DEVICE
Sets the device that will be used in a target region. The OpenMP routine
-``omp_set_default_device`` or a device clause in a parallel pragma can override
+`omp_set_default_device` or a device clause in a parallel pragma can override
this variable. If no device with the specified device number exists, the code is
executed on the host. If this environment variable is not set, device number 0
is used.
-OMP_DYNAMIC
-"""""""""""
+#### OMP_DYNAMIC
-Enables (``true``) or disables (``false``) the dynamic adjustment of the
+Enables (`true`) or disables (`false`) the dynamic adjustment of the
number of threads.
-| **Default:** ``false``
+**Default:**
+
+
-OMP_MAX_ACTIVE_LEVELS
-"""""""""""""""""""""
+`false`
+
+#### OMP_MAX_ACTIVE_LEVELS
The maximum number of levels of parallel nesting for the program.
-| **Default:** ``1``
+**Default:**
-OMP_NESTED
-""""""""""
+
-.. warning::
- Deprecated. Please use ``OMP_MAX_ACTIVE_LEVELS`` to control nested parallelism
+`1`
-Enables (``true``) or disables (``false``) nested parallelism.
+#### OMP_NESTED
+
+:::{warning}
+Deprecated. Please use `OMP_MAX_ACTIVE_LEVELS` to control nested parallelism
+:::
+
+Enables (`true`) or disables (`false`) nested parallelism.
+
+**Default:**
-| **Default:** ``false``
+
-OMP_NUM_THREADS
-"""""""""""""""
+`false`
+
+#### OMP_NUM_THREADS
Sets the maximum number of threads to use for OpenMP parallel regions if no
other value is specified in the application.
@@ -111,16 +112,27 @@ list is left out, it implies the normal default value for threads is used at the
outer-most level. If the integer is left out of any other level, the number of
threads for that level is inherited from the previous level.
-| **Default:** The number of processors visible to the operating system on which the program is executed.
-| **Syntax:** ``OMP_NUM_THREADS=value[,value]*``
-| **Example:** ``OMP_NUM_THREADS=4,3``
+**Default:**
+
+ The number of processors visible to the operating system on which the program is executed.
+
+**Syntax:**
+
+
+
+`OMP_NUM_THREADS=value[,value]*`
+
+**Example:**
+
+
-OMP_PLACES
-""""""""""
+`OMP_NUM_THREADS=4,3`
+
+#### OMP_PLACES
Specifies an explicit ordered list of places, either as an abstract name
describing a set of places or as an explicit list of places described by
-non-negative numbers. An exclusion operator, ``!``, can also be used to exclude
+non-negative numbers. An exclusion operator, `!`, can also be used to exclude
the number or place immediately following the operator.
For **explicit lists**, an ordered list of places is specified with each place
@@ -130,132 +142,168 @@ operating system logical processor numbers and can be thought of as an OS affini
Individual places can be specified through two methods.
Both the **examples** below represent the same place.
-* An explicit list of comma-separated non-negatives numbers **Example:** ``{0,2,4,6}``
-* An interval with notation ``<lower-bound>:<length>[:<stride>]``. **Example:** ``{0:4:2}``. When ``<stride>`` is omitted, a unit stride is assumed.
+- An explicit list of comma-separated non-negatives numbers **Example:** `{0,2,4,6}`
+- An interval with notation `<lower-bound>:<length>[:<stride>]`. **Example:** `{0:4:2}`. When `<stride>` is omitted, a unit stride is assumed.
The interval notation represents this set of numbers:
-::
-
- <lower-bound>, <lower-bound> + <stride>, ..., <lower-bound> + (<length> - 1) * <stride>
-
+```
+<lower-bound>, <lower-bound> + <stride>, ..., <lower-bound> + (<length> - 1) * <stride>
+```
A place list can also be specified using the same interval
-notation: ``{place}:<length>[:<stride>]``.
-This represents the list of length ``<length>`` places determined by the following:
-
-.. code-block:: c
+notation: `{place}:<length>[:<stride>]`.
+This represents the list of length `<length>` places determined by the following:
- {place}, {place} + <stride>, ..., {place} + (<length>-1)*<stride>
- Where given {place} and integer N, {place} + N = {place with every number offset by N}
- Example: {0,3,6}:4:1 represents {0,3,6}, {1,4,7}, {2,5,8}, {3,6,9}
+```c
+{place}, {place} + <stride>, ..., {place} + (<length>-1)*<stride>
+Where given {place} and integer N, {place} + N = {place with every number offset by N}
+Example: {0,3,6}:4:1 represents {0,3,6}, {1,4,7}, {2,5,8}, {3,6,9}
+```
**Examples of explicit lists:**
These all represent the same set of places
-::
+```
+OMP_PLACES="{0,1,2,3},{4,5,6,7},{8,9,10,11},{12,13,14,15}"
+OMP_PLACES="{0:4},{4:4},{8:4},{12:4}"
+OMP_PLACES="{0:4}:4:4"
+```
- OMP_PLACES="{0,1,2,3},{4,5,6,7},{8,9,10,11},{12,13,14,15}"
- OMP_PLACES="{0:4},{4:4},{8:4},{12:4}"
- OMP_PLACES="{0:4}:4:4"
-
-.. note::
- When specifying a place using a set of numbers, if any number cannot be
- mapped to a processor on the target platform, then that number is
- ignored within the place, but the rest of the place is kept intact.
- If all numbers within a place are invalid, then the entire place is removed
- from the place list, but the rest of place list is kept intact.
+:::{note}
+When specifying a place using a set of numbers, if any number cannot be
+mapped to a processor on the target platform, then that number is
+ignored within the place, but the rest of the place is kept intact.
+If all numbers within a place are invalid, then the entire place is removed
+from the place list, but the rest of place list is kept intact.
+:::
The **abstract names** listed below are understood by the run-time environment:
-* ``threads:`` Each place corresponds to a single hardware thread.
-* ``cores:`` Each place corresponds to a single core (having one or more hardware threads).
-* ``sockets:`` Each place corresponds to a single socket (consisting of one or more cores).
-* ``numa_domains:`` Each place corresponds to a single NUMA domain (consisting of one or more cores).
-* ``ll_caches:`` Each place corresponds to a last-level cache (consisting of one or more cores).
+- `threads:` Each place corresponds to a single hardware thread.
+- `cores:` Each place corresponds to a single core (having one or more hardware threads).
+- `sockets:` Each place corresponds to a single socket (consisting of one or more cores).
+- `numa_domains:` Each place corresponds to a single NUMA domain (consisting of one or more cores).
+- `ll_caches:` Each place corresponds to a last-level cache (consisting of one or more cores).
The abstract name may be appended by a positive number in parentheses to
-denote the length of the place list to be created, that is ``abstract_name(num-places)``.
+denote the length of the place list to be created, that is `abstract_name(num-places)`.
If the optional number isn't specified, then the runtime will use all available
-resources of type ``abstract_name``. When requesting fewer places than available
-on the system, the first available resources as determined by ``abstract_name``
+resources of type `abstract_name`. When requesting fewer places than available
+on the system, the first available resources as determined by `abstract_name`
are used. When requesting more places than available on the system, only the
available resources are used.
**Examples of abstract names:**
-::
- OMP_PLACES=threads
- OMP_PLACES=threads(4)
+```
+OMP_PLACES=threads
+OMP_PLACES=threads(4)
+```
+
+#### OMP_PROC_BIND (Windows, Linux)
-OMP_PROC_BIND (Windows, Linux)
-""""""""""""""""""""""""""""""
Sets the thread affinity policy to be used for parallel regions at the
-corresponding nested level. Enables (``true``) or disables (``false``)
+corresponding nested level. Enables (`true`) or disables (`false`)
the binding of threads to processor contexts. If enabled, this is the
-same as specifying ``KMP_AFFINITY=scatter``. If disabled, this is the
-same as specifying ``KMP_AFFINITY=none``.
+same as specifying `KMP_AFFINITY=scatter`. If disabled, this is the
+same as specifying `KMP_AFFINITY=none`.
-**Acceptable values:** ``true``, ``false``, or a comma separated list, each
-element of which is one of the following values: ``master``, ``close``, ``spread``, or ``primary``.
+**Acceptable values:** `true`, `false`, or a comma separated list, each
+element of which is one of the following values: `master`, `close`, `spread`, or `primary`.
-**Default:** ``false``
+**Default:** `false`
-.. warning::
- ``master`` is deprecated. The semantics of ``master`` are the same as ``primary``.
+:::{warning}
+`master` is deprecated. The semantics of `master` are the same as `primary`.
+:::
-If set to ``false``, the execution environment may move OpenMP threads between
-OpenMP places, thread affinity is disabled, and ``proc_bind`` clauses on
+If set to `false`, the execution environment may move OpenMP threads between
+OpenMP places, thread affinity is disabled, and `proc_bind` clauses on
parallel constructs are ignored. Otherwise, the execution environment should
not move OpenMP threads between OpenMP places, thread affinity is enabled, and
the initial thread is bound to the first place in the OpenMP place list.
-If set to ``primary``, all threads are bound to the same place as the primary
+If set to `primary`, all threads are bound to the same place as the primary
thread.
-If set to ``close``, threads are bound to successive places, near where the
+If set to `close`, threads are bound to successive places, near where the
primary thread is bound.
-If set to ``spread``, the primary thread's partition is subdivided and threads
+If set to `spread`, the primary thread's partition is subdivided and threads
are bound to single place successive sub-partitions.
-| **Related environment variables:** ``KMP_AFFINITY`` (overrides ``OMP_PROC_BIND``).
+**Related environment variables:**
+
+
+
+`KMP_AFFINITY`
+
+ (overrides
+
+`OMP_PROC_BIND`
+
+).
+
+#### OMP_SCHEDULE
-OMP_SCHEDULE
-""""""""""""
Sets the run-time schedule type and an optional chunk size.
-| **Default:** ``static``, no chunk size specified
-| **Syntax:** ``OMP_SCHEDULE="kind[,chunk_size]"``
+**Default:**
+
+
+
+`static`
+
+, no chunk size specified
-OMP_STACKSIZE
-"""""""""""""
+**Syntax:**
+
+
+
+`OMP_SCHEDULE="kind[,chunk_size]"`
+
+#### OMP_STACKSIZE
Sets the number of bytes to allocate for each OpenMP thread to use as the
private stack for the thread. Recommended size is 16M.
-Use the optional suffixes to specify byte units: ``B`` (bytes), ``K`` (Kilobytes),
-``M`` (Megabytes), ``G`` (Gigabytes), or ``T`` (Terabytes) to specify the units.
+Use the optional suffixes to specify byte units: `B` (bytes), `K` (Kilobytes),
+`M` (Megabytes), `G` (Gigabytes), or `T` (Terabytes) to specify the units.
If you specify a value without a suffix, the byte unit
-is assumed to be ``K`` (Kilobytes).
+is assumed to be `K` (Kilobytes).
This variable does not affect the native operating system threads created by the
user program, or the thread executing the sequential part of an OpenMP program.
-The ``kmp_{set,get}_stacksize_s()`` routines set/retrieve the value.
-The ``kmp_set_stacksize_s()`` routine must be called from sequential part, before
-first parallel region is created. Otherwise, calling ``kmp_set_stacksize_s()``
+The `kmp_{set,get}_stacksize_s()` routines set/retrieve the value.
+The `kmp_set_stacksize_s()` routine must be called from sequential part, before
+first parallel region is created. Otherwise, calling `kmp_set_stacksize_s()`
has no effect.
-| **Default:**
+**Default:**
+
+- 32-bit architecture: `2M`
+- 64-bit architecture: `4M`
+
+**Related environment variables:**
+
+
+
+`KMP_STACKSIZE`
+
+ (overrides
-* 32-bit architecture: ``2M``
-* 64-bit architecture: ``4M``
+`OMP_STACKSIZE`
-| **Related environment variables:** ``KMP_STACKSIZE`` (overrides ``OMP_STACKSIZE``).
-| **Example:** ``OMP_STACKSIZE=8M``
+).
-OMP_THREAD_LIMIT
-""""""""""""""""
+**Example:**
+
+
+
+`OMP_STACKSIZE=8M`
+
+#### OMP_THREAD_LIMIT
Limits the number of simultaneously-executing threads in an OpenMP program.
@@ -265,136 +313,164 @@ If this limit is reached when an OpenMP parallel region begins, a one-time
warning message might be generated indicating that the number of threads in
the team was reduced, but the program will continue.
-The ``omp_get_thread_limit()`` routine returns the value of the limit.
+The `omp_get_thread_limit()` routine returns the value of the limit.
+
+**Default:**
+
+ No enforced limit
+
+**Related environment variable:**
+
+
-| **Default:** No enforced limit
-| **Related environment variable:** ``KMP_ALL_THREADS`` (overrides ``OMP_THREAD_LIMIT``).
+`KMP_ALL_THREADS`
-OMP_WAIT_POLICY
-"""""""""""""""
+ (overrides
+
+`OMP_THREAD_LIMIT`
+
+).
+
+#### OMP_WAIT_POLICY
Decides whether threads spin (active) or yield (passive) while they are waiting.
-``OMP_WAIT_POLICY=active`` is an alias for ``KMP_LIBRARY=turnaround``, and
-``OMP_WAIT_POLICY=passive`` is an alias for ``KMP_LIBRARY=throughput``.
+`OMP_WAIT_POLICY=active` is an alias for `KMP_LIBRARY=turnaround`, and
+`OMP_WAIT_POLICY=passive` is an alias for `KMP_LIBRARY=throughput`.
+
+**Default:**
+
+
-| **Default:** ``passive``
+`passive`
-.. note::
- Although the default is ``passive``, unless the user has explicitly set
- ``OMP_WAIT_POLICY``, there is a small period of active spinning determined
- by ``KMP_BLOCKTIME``.
+:::{note}
+Although the default is `passive`, unless the user has explicitly set
+`OMP_WAIT_POLICY`, there is a small period of active spinning determined
+by `KMP_BLOCKTIME`.
+:::
-KMP_AFFINITY (Windows, Linux)
-"""""""""""""""""""""""""""""
+#### KMP_AFFINITY (Windows, Linux)
Enables run-time library to bind threads to physical processing units.
You must set this environment variable before the first parallel region, or
-certain API calls including ``omp_get_max_threads()``, ``omp_get_num_procs()``
+certain API calls including `omp_get_max_threads()`, `omp_get_num_procs()`
and any affinity API calls.
-**Syntax:** ``KMP_AFFINITY=[<modifier>,...]<type>[,<permute>][,<offset>]``
+**Syntax:** `KMP_AFFINITY=[<modifier>,...]<type>[,<permute>][,<offset>]`
-``modifiers`` are optional strings consisting of a keyword and possibly a specifier
+`modifiers` are optional strings consisting of a keyword and possibly a specifier
-* ``respect`` (default) and ``norespect`` - determine whether to respect the original process affinity mask.
-* ``verbose`` and ``noverbose`` (default) - determine whether to display affinity information.
-* ``warnings`` (default) and ``nowarnings`` - determine whether to display warnings during affinity detection.
-* ``reset`` and ``noreset`` (default) - determine whether to reset primary thread's affinity after outermost parallel region(s)
-* ``granularity=<specifier>`` - takes the following specifiers ``thread``, ``core`` (default), ``tile``,
- ``socket``, ``die``, ``group`` (Windows only).
+- `respect` (default) and `norespect` - determine whether to respect the original process affinity mask.
+- `verbose` and `noverbose` (default) - determine whether to display affinity information.
+- `warnings` (default) and `nowarnings` - determine whether to display warnings during affinity detection.
+- `reset` and `noreset` (default) - determine whether to reset primary thread's affinity after outermost parallel region(s)
+- `granularity=<specifier>` - takes the following specifiers `thread`, `core` (default), `tile`,
+ `socket`, `die`, `group` (Windows only).
The granularity describes the lowest topology levels that OpenMP threads are allowed to float within a topology map.
- For example, if ``granularity=core``, then the OpenMP threads will be allowed to move between logical processors within
- a single core. If ``granularity=thread``, then the OpenMP threads will be restricted to a single logical processor.
-* ``proclist=[<proc_list>]`` - The ``proc_list`` is specified by
-
-+--------------------+----------------------------------------+
-| Value | Description |
-+====================+========================================+
-| <proc_list> := | <proc_id> | { <id_list> } |
-+--------------------+----------------------------------------+
-| <id_list> := | <proc_id> | <proc_id>,<id_list> |
-+--------------------+----------------------------------------+
-
-Where each ``proc_id`` represents an operating system logical processor ID.
-For example, ``proclist=[3,0,{1,2},{0,3}]`` with ``OMP_NUM_THREADS=4`` would place thread 0 on
+ For example, if `granularity=core`, then the OpenMP threads will be allowed to move between logical processors within
+ a single core. If `granularity=thread`, then the OpenMP threads will be restricted to a single logical processor.
+- `proclist=[<proc_list>]` - The `proc_list` is specified by
+
+| Value | Description |
+| --------------- | ----------------------------------- |
+| \<proc_list> := | \<proc_id> \| { \<id_list> } |
+| \<id_list> := | \<proc_id> \| \<proc_id>,\<id_list> |
+
+Where each `proc_id` represents an operating system logical processor ID.
+For example, `proclist=[3,0,{1,2},{0,3}]` with `OMP_NUM_THREADS=4` would place thread 0 on
OS logical processor 3, thread 1 on OS logical processor 0, thread 2 on both OS logical
processors 1 & 2, and thread 3 on OS logical processors 0 & 3.
-``type`` is the thread affinity policy to choose.
-Valid choices are ``none``, ``balanced``, ``compact``, ``scatter``, ``explicit``, ``disabled``
+`type` is the thread affinity policy to choose.
+Valid choices are `none`, `balanced`, `compact`, `scatter`, `explicit`, `disabled`
-* type ``none`` (default) - Does not bind OpenMP threads to particular thread contexts;
+- type `none` (default) - Does not bind OpenMP threads to particular thread contexts;
however, if the operating system supports affinity, the compiler still uses the
OpenMP thread affinity interface to determine machine topology.
- Specify ``KMP_AFFINITY=verbose,none`` to list a machine topology map.
-* type ``compact`` - Specifying compact assigns the OpenMP thread <n>+1 to a free thread
- context as close as possible to the thread context where the <n> OpenMP thread was
+ Specify `KMP_AFFINITY=verbose,none` to list a machine topology map.
+- type `compact` - Specifying compact assigns the OpenMP thread \<n>+1 to a free thread
+ context as close as possible to the thread context where the \<n> OpenMP thread was
placed. For example, in a topology map, the nearer a node is to the root, the more
significance the node has when sorting the threads.
-* type ``scatter`` - Specifying scatter distributes the threads as evenly as
- possible across the entire system. ``scatter`` is the opposite of ``compact``; so the
+- type `scatter` - Specifying scatter distributes the threads as evenly as
+ possible across the entire system. `scatter` is the opposite of `compact`; so the
leaves of the node are most significant when sorting through the machine topology map.
-* type ``balanced`` - Places threads on separate cores until all cores have at least one thread,
- similar to the ``scatter`` type. However, when the runtime must use multiple hardware thread
+- type `balanced` - Places threads on separate cores until all cores have at least one thread,
+ similar to the `scatter` type. However, when the runtime must use multiple hardware thread
contexts on the same core, the balanced type ensures that the OpenMP thread numbers are close
to each other, which scatter does not do. This affinity type is supported on the CPU only for
single socket systems.
-* type ``explicit`` - Specifying explicit assigns OpenMP threads to a list of OS proc IDs that
- have been explicitly specified by using the ``proclist`` modifier, which is required
+- type `explicit` - Specifying explicit assigns OpenMP threads to a list of OS proc IDs that
+ have been explicitly specified by using the `proclist` modifier, which is required
for this affinity type.
-* type ``disabled`` - Specifying disabled completely disables the thread affinity interfaces.
+- type `disabled` - Specifying disabled completely disables the thread affinity interfaces.
This forces the OpenMP run-time library to behave as if the affinity interface was not
supported by the operating system. This includes the low-level API interfaces such
- as ``kmp_set_affinity`` and ``kmp_get_affinity``, which have no effect and will return
+ as `kmp_set_affinity` and `kmp_get_affinity`, which have no effect and will return
a nonzero error code.
-For both ``compact`` and ``scatter``, ``permute`` and ``offset`` are allowed;
+For both `compact` and `scatter`, `permute` and `offset` are allowed;
however, if you specify only one integer, the runtime interprets the value as
a permute specifier. **Both permute and offset default to 0.**
-The ``permute`` specifier controls which levels are most significant when sorting
-the machine topology map. A value for ``permute`` forces the mappings to make the
+The `permute` specifier controls which levels are most significant when sorting
+the machine topology map. A value for `permute` forces the mappings to make the
specified number of most significant levels of the sort the least significant,
and it inverts the order of significance. The root node of the tree is not
considered a separate level for the sort operations.
-The ``offset`` specifier indicates the starting position for thread assignment.
+The `offset` specifier indicates the starting position for thread assignment.
+
+**Default:**
+
+
+
+`noverbose,warnings,respect,granularity=core,none`
+
+**Related environment variable:**
+
+
+
+`OMP_PROC_BIND`
+
+ (
-| **Default:** ``noverbose,warnings,respect,granularity=core,none``
-| **Related environment variable:** ``OMP_PROC_BIND`` (``KMP_AFFINITY`` takes precedence)
+`KMP_AFFINITY`
-.. note::
- On Windows with multiple processor groups, the norespect affinity modifier
- is assumed when the process affinity mask equals a single processor group
- (which is default on Windows). Otherwise, the respect affinity modifier is used.
+ takes precedence)
-.. note::
- On Windows with multiple processor groups, if the granularity is too coarse, it
- will be set to ``granularity=group``. For example, if two processor groups exist
- across one socket, and ``granularity=socket`` the runtime will shift the
- granularity down to group since that is the largest granularity allowed by the OS.
+:::{note}
+On Windows with multiple processor groups, the norespect affinity modifier
+is assumed when the process affinity mask equals a single processor group
+(which is default on Windows). Otherwise, the respect affinity modifier is used.
+:::
-KMP_HIDDEN_HELPER_AFFINITY (Windows, Linux)
-"""""""""""""""""""""""""""""""""""""""""""
+:::{note}
+On Windows with multiple processor groups, if the granularity is too coarse, it
+will be set to `granularity=group`. For example, if two processor groups exist
+across one socket, and `granularity=socket` the runtime will shift the
+granularity down to group since that is the largest granularity allowed by the OS.
+:::
+
+#### KMP_HIDDEN_HELPER_AFFINITY (Windows, Linux)
Enables run-time library to bind hidden helper threads to physical processing units.
-This environment variable has the same syntax and semantics as ``KMP_AFFINIY`` but only
+This environment variable has the same syntax and semantics as `KMP_AFFINIY` but only
applies to the hidden helper team.
You must set this environment variable before the first parallel region, or
-certain API calls including ``omp_get_max_threads()``, ``omp_get_num_procs()``
+certain API calls including `omp_get_max_threads()`, `omp_get_num_procs()`
and any affinity API calls.
-**Syntax:** Same as ``KMP_AFFINITY``
+**Syntax:** Same as `KMP_AFFINITY`
-The following ``modifiers`` are ignored in ``KMP_HIDDEN_HELPER_AFFINITY`` and are only valid
-for ``KMP_AFFINITY``:
-* ``respect`` and ``norespect``
-* ``reset`` and ``noreset``
+The following `modifiers` are ignored in `KMP_HIDDEN_HELPER_AFFINITY` and are only valid
+for `KMP_AFFINITY`:
+\* `respect` and `norespect`
+\* `reset` and `noreset`
-KMP_ALL_THREADS
-"""""""""""""""
+#### KMP_ALL_THREADS
Limits the number of simultaneously-executing threads in an OpenMP program.
If this limit is reached and another native operating system thread encounters
@@ -403,85 +479,114 @@ message. If this limit is reached at the time an OpenMP parallel region begins,
a one-time warning message may be generated indicating that the number of
threads in the team was reduced, but the program will continue execution.
-| **Default:** No enforced limit.
-| **Related environment variable:** ``OMP_THREAD_LIMIT`` (``KMP_ALL_THREADS`` takes precedence)
+**Default:**
+
+ No enforced limit.
+
+**Related environment variable:**
+
+
+
+`OMP_THREAD_LIMIT`
+
+ (
+
+`KMP_ALL_THREADS`
+
+ takes precedence)
-KMP_BLOCKTIME
-"""""""""""""
+#### KMP_BLOCKTIME
Sets the time that a thread should wait, after completing the
execution of a parallel region, before sleeping.
-Use the optional suffixes: ``ms`` (milliseconds), or ``us`` (microseconds) to
+Use the optional suffixes: `ms` (milliseconds), or `us` (microseconds) to
specify/change the units. Defaults units is milliseconds.
-Specify ``infinite`` for an unlimited wait time.
+Specify `infinite` for an unlimited wait time.
+
+**Default:**
+
+ 200 milliseconds
+
+**Related Environment Variable:**
+
+
+
+`KMP_LIBRARY`
+
+**Example:**
+
+
-| **Default:** 200 milliseconds
-| **Related Environment Variable:** ``KMP_LIBRARY``
-| **Example:** ``KMP_BLOCKTIME=1ms``
+`KMP_BLOCKTIME=1ms`
-KMP_CPUINFO_FILE
-""""""""""""""""
+#### KMP_CPUINFO_FILE
Specifies an alternate file name for a file containing the machine topology
-description. The file must be in the same format as :file:`/proc/cpuinfo`.
+description. The file must be in the same format as {file}`/proc/cpuinfo`.
**Default:** None
-KMP_DETERMINISTIC_REDUCTION
-"""""""""""""""""""""""""""
+#### KMP_DETERMINISTIC_REDUCTION
-Enables (``true``) or disables (``false``) the use of a specific ordering of
+Enables (`true`) or disables (`false`) the use of a specific ordering of
the reduction operations for implementing the reduction clause for an OpenMP
parallel region. This has the effect that, for a given number of threads, in
a given parallel region, for a given data set and reduction operation, a
floating point reduction done for an OpenMP reduction clause has a consistent
floating point result from run to run, since round-off errors are identical.
-| **Default:** ``false``
-| **Example:** ``KMP_DETERMINISTIC_REDUCTION=true``
+**Default:**
+
+
+
+`false`
+
+**Example:**
-KMP_DYNAMIC_MODE
-""""""""""""""""
+
+
+`KMP_DETERMINISTIC_REDUCTION=true`
+
+#### KMP_DYNAMIC_MODE
Selects the method used to determine the number of threads to use for a parallel
-region when ``OMP_DYNAMIC=true``. Possible values: (``load_balance`` | ``thread_limit``), where,
+region when `OMP_DYNAMIC=true`. Possible values: (`load_balance` | `thread_limit`), where,
-* ``load_balance``: tries to avoid using more threads than available execution units on the machine;
-* ``thread_limit``: tries to avoid using more threads than total execution units on the machine.
+- `load_balance`: tries to avoid using more threads than available execution units on the machine;
+- `thread_limit`: tries to avoid using more threads than total execution units on the machine.
-**Default:** ``load_balance`` (on all supported platforms)
+**Default:** `load_balance` (on all supported platforms)
+
+#### KMP_HOT_TEAMS_MAX_LEVEL
-KMP_HOT_TEAMS_MAX_LEVEL
-"""""""""""""""""""""""
Sets the maximum nested level to which teams of threads will be hot.
-.. note::
- A hot team is a team of threads optimized for faster reuse by subsequent
- parallel regions. In a hot team, threads are kept ready for execution of
- the next parallel region, in contrast to the cold team, which is freed
- after each parallel region, with its threads going into a common pool
- of threads.
+:::{note}
+A hot team is a team of threads optimized for faster reuse by subsequent
+parallel regions. In a hot team, threads are kept ready for execution of
+the next parallel region, in contrast to the cold team, which is freed
+after each parallel region, with its threads going into a common pool
+of threads.
+:::
For values of 2 and above, nested parallelism should be enabled.
**Default:** 1
-KMP_HOT_TEAMS_MODE
-""""""""""""""""""
+#### KMP_HOT_TEAMS_MODE
Specifies the run-time behavior when the number of threads in a hot team is reduced.
Possible values:
-* ``0`` - Extra threads are freed and put into a common pool of threads.
-* ``1`` - Extra threads are kept in the team in reserve, for faster reuse
+- `0` - Extra threads are freed and put into a common pool of threads.
+- `1` - Extra threads are kept in the team in reserve, for faster reuse
in subsequent parallel regions.
**Default:** 0
-KMP_HW_SUBSET
-"""""""""""""
+#### KMP_HW_SUBSET
Specifies the subset of available hardware resources for the hardware topology
hierarchy. The subset is specified in terms of number of units per upper layer
@@ -492,61 +597,83 @@ or a limiting process affinity mask. You can also specify an offset value to set
which resources to use. When available, you can specify attributes to select
different subsets of resources.
-An extended syntax is available when ``KMP_TOPOLOGY_METHOD=hwloc``. Depending on what
+An extended syntax is available when `KMP_TOPOLOGY_METHOD=hwloc`. Depending on what
resources are detected, you may be able to specify additional resources, such as
NUMA domains and groups of hardware resources that share certain cache levels.
-**Basic syntax:** ``[:][num_units|*]ID[@offset][:attribute] [,[num_units|*]ID[@offset][:attribute]...]``
+**Basic syntax:** `[:][num_units|*]ID[@offset][:attribute] [,[num_units|*]ID[@offset][:attribute]...]`
An optional colon (:) can be specified at the beginning of the syntax to specify an explicit hardware subset. The default is an implicit hardware subset.
Supported unit IDs are not case-insensitive.
-| ``S`` - socket
-| ``num_units`` specifies the requested number of sockets.
+`S`
+
+ - socket
+
+`num_units`
+
+ specifies the requested number of sockets.
+
+`D`
+
+ - die
+
+`num_units`
+
+ specifies the requested number of dies per socket.
+
+`C`
+
+ - core
-| ``D`` - die
-| ``num_units`` specifies the requested number of dies per socket.
+`num_units`
-| ``C`` - core
-| ``num_units`` specifies the requested number of cores per die - if any - otherwise, per socket.
+ specifies the requested number of cores per die - if any - otherwise, per socket.
-| ``T`` - thread
-| ``num_units`` specifies the requested number of HW threads per core.
+`T`
-.. note::
- ``num_units`` can be left out or explicitly specified as ``*`` instead of a positive integer
- meaning use all specified resources at that level.
- e.g., ``1s,*c`` means use 1 socket and all the cores on that socket
+ - thread
-``offset`` - (Optional) The number of units to skip.
+`num_units`
-``attribute`` - (Optional) An attribute differentiating resources at a particular level. The attributes available to users are:
+ specifies the requested number of HW threads per core.
-* **Core type** - On Intel architectures, this can be ``intel_atom`` or ``intel_core``
-* **Core efficiency** - This is specified as ``eff``:emphasis:`num` where :emphasis:`num` is a number from 0
+:::{note}
+`num_units` can be left out or explicitly specified as `*` instead of a positive integer
+meaning use all specified resources at that level.
+e.g., `1s,*c` means use 1 socket and all the cores on that socket
+:::
+
+`offset` - (Optional) The number of units to skip.
+
+`attribute` - (Optional) An attribute differentiating resources at a particular level. The attributes available to users are:
+
+- **Core type** - On Intel architectures, this can be `intel_atom` or `intel_core`
+- **Core efficiency** - This is specified as `eff`{emphasis}`num` where {emphasis}`num` is a number from 0
to the number of core efficiencies detected in the machine topology minus one.
- E.g., ``eff0``. The greater the efficiency number the more performant the core. There may be
- more core efficiencies than core types and can be viewed by setting ``KMP_AFFINITY=verbose``
+ E.g., `eff0`. The greater the efficiency number the more performant the core. There may be
+ more core efficiencies than core types and can be viewed by setting `KMP_AFFINITY=verbose`
-.. note::
- The hardware cache can be specified as a unit, e.g. L2 for L2 cache,
- or LL for last level cache.
+:::{note}
+The hardware cache can be specified as a unit, e.g. L2 for L2 cache,
+or LL for last level cache.
+:::
**Extended syntax when KMP_TOPOLOGY_METHOD=hwloc:**
Additional IDs can be specified if detected. For example:
-``N`` - numa
-``num_units`` specifies the requested number of NUMA nodes per upper layer
+`N` - numa
+`num_units` specifies the requested number of NUMA nodes per upper layer
unit, e.g. per socket.
-``TI`` - tile
+`TI` - tile
num_units specifies the requested number of tiles to use per upper layer
unit, e.g. per NUMA node.
-When any numa or tile units are specified in ``KMP_HW_SUBSET`` and the hwloc
-topology method is available, the ``KMP_TOPOLOGY_METHOD`` will be automatically
+When any numa or tile units are specified in `KMP_HW_SUBSET` and the hwloc
+topology method is available, the `KMP_TOPOLOGY_METHOD` will be automatically
set to hwloc, so there is no need to set it explicitly.
For an **explicit hardware subset**, if one or more topology layers detected by the
@@ -565,239 +692,233 @@ If you don't specify one or more types of resource, such as socket or thread,
all available resources of that type are used.
The run-time library prints a warning, and the setting of
-``KMP_HW_SUBSET`` is ignored if:
+`KMP_HW_SUBSET` is ignored if:
-* a resource is specified, but detection of that resource is not supported
+- a resource is specified, but detection of that resource is not supported
by the chosen topology detection method and/or
-* a resource is specified twice. An exception to this condition is if attributes
+- a resource is specified twice. An exception to this condition is if attributes
differentiate the resource.
-* attributes are used when not detected in the machine topology or conflict with
+- attributes are used when not detected in the machine topology or conflict with
each other.
-This variable does not work if ``KMP_AFFINITY=disabled``.
+This variable does not work if `KMP_AFFINITY=disabled`.
**Default:** If omitted, the default value is to use all the
available hardware resources.
**Implicit Hardware Subset Examples:**
-* ``2s,4c,2t``: Use the first 2 sockets (s0 and s1), the first 4 cores on each
+- `2s,4c,2t`: Use the first 2 sockets (s0 and s1), the first 4 cores on each
socket (c0 - c3), and 2 threads per core.
-* ``2s at 2,4c at 8,2t``: Skip the first 2 sockets (s0 and s1) and use 2 sockets
+- `2s at 2,4c at 8,2t`: Skip the first 2 sockets (s0 and s1) and use 2 sockets
(s2-s3), skip the first 8 cores (c0-c7) and use 4 cores on each socket
(c8-c11), and use 2 threads per core.
-* ``5C at 1,3T``: Use all available sockets, skip the first core and use 5 cores,
+- `5C at 1,3T`: Use all available sockets, skip the first core and use 5 cores,
and use 3 threads per core.
-* ``1T``: Use all cores on all sockets, 1 thread per core.
-* ``1s, 1d, 1n, 1c, 1t``: Use 1 socket, 1 die, 1 NUMA node, 1 core, 1 thread
- - use HW thread as a result.
-* ``4c:intel_atom,5c:intel_core``: Use all available sockets and use 4
+- `1T`: Use all cores on all sockets, 1 thread per core.
+- `1s, 1d, 1n, 1c, 1t`: Use 1 socket, 1 die, 1 NUMA node, 1 core, 1 thread
+ \- use HW thread as a result.
+- `4c:intel_atom,5c:intel_core`: Use all available sockets and use 4
Intel Atom(R) processor cores and 5 Intel(R) Core(TM) processor cores per socket.
-* ``2c:eff0 at 1,3c:eff1``: Use all available sockets, skip the first core with efficiency 0
+- `2c:eff0 at 1,3c:eff1`: Use all available sockets, skip the first core with efficiency 0
and use the next 2 cores with efficiency 0 and 3 cores with efficiency 1 per socket.
-* ``1s, 1c, 1t``: Use 1 socket, 1 core, 1 thread. This may result in using
+- `1s, 1c, 1t`: Use 1 socket, 1 core, 1 thread. This may result in using
single thread on a 3-layer topology architecture, or multiple threads on
4-layer or 5-layer architecture. Result may even be different on the same
- architecture, depending on ``KMP_TOPOLOGY_METHOD`` specified, as hwloc can
+ architecture, depending on `KMP_TOPOLOGY_METHOD` specified, as hwloc can
often detect more topology layers than the default method used by the OpenMP
run-time library.
-* ``*c:eff1 at 3``: Use all available sockets, skip the first three cores of
+- `*c:eff1 at 3`: Use all available sockets, skip the first three cores of
efficiency 1, and then use the rest of the available cores of efficiency 1.
Explicit Hardware Subset Examples:
-* ``:2s,6t`` Use exactly the first two sockets and 6 threads per socket.
-* ``:1t at 7`` Skip the first 7 threads (t0-t6) and use exactly one thread (t7).
-* ``:5c,1t`` Use exactly the first 5 cores (c0-c4) and the first thread on each core.
+- `:2s,6t` Use exactly the first two sockets and 6 threads per socket.
+- `:1t at 7` Skip the first 7 threads (t0-t6) and use exactly one thread (t7).
+- `:5c,1t` Use exactly the first 5 cores (c0-c4) and the first thread on each core.
-To see the result of the setting, you can specify ``verbose`` modifier in
-``KMP_AFFINITY`` environment variable. The OpenMP run-time library will output
-to ``stderr`` the information about the discovered hardware topology before and
-after the ``KMP_HW_SUBSET`` setting was applied.
+To see the result of the setting, you can specify `verbose` modifier in
+`KMP_AFFINITY` environment variable. The OpenMP run-time library will output
+to `stderr` the information about the discovered hardware topology before and
+after the `KMP_HW_SUBSET` setting was applied.
-KMP_INHERIT_FP_CONTROL
-""""""""""""""""""""""
+#### KMP_INHERIT_FP_CONTROL
-Enables (``true``) or disables (``false``) the copying of the floating-point
+Enables (`true`) or disables (`false`) the copying of the floating-point
control settings of the primary thread to the floating-point control settings
of the OpenMP worker threads at the start of each parallel region.
-**Default:** ``true``
+**Default:** `true`
-KMP_LIBRARY
-"""""""""""
+#### KMP_LIBRARY
Selects the OpenMP run-time library execution mode. The values for this variable
-are ``serial``, ``turnaround``, or ``throughput``.
+are `serial`, `turnaround`, or `throughput`.
+
+**Default:**
+
+
+
+`throughput`
+
+**Related environment variable:**
-| **Default:** ``throughput``
-| **Related environment variable:** ``KMP_BLOCKTIME`` and ``OMP_WAIT_POLICY``
+
-KMP_SETTINGS
-""""""""""""
+`KMP_BLOCKTIME`
-Enables (``true``) or disables (``false``) the printing of OpenMP run-time library
+ and
+
+`OMP_WAIT_POLICY`
+
+#### KMP_SETTINGS
+
+Enables (`true`) or disables (`false`) the printing of OpenMP run-time library
environment variables during program execution. Two lists of variables are printed:
user-defined environment variables settings and effective values of variables used
by OpenMP run-time library.
-**Default:** ``false``
+**Default:** `false`
-KMP_STACKSIZE
-"""""""""""""
+#### KMP_STACKSIZE
Sets the number of bytes to allocate for each OpenMP thread to use as its private stack.
-Recommended size is ``16M``.
+Recommended size is `16M`.
-Use the optional suffixes to specify byte units: ``B`` (bytes), ``K`` (Kilobytes),
-``M`` (Megabytes), ``G`` (Gigabytes), or ``T`` (Terabytes) to specify the units.
+Use the optional suffixes to specify byte units: `B` (bytes), `K` (Kilobytes),
+`M` (Megabytes), `G` (Gigabytes), or `T` (Terabytes) to specify the units.
If you specify a value without a suffix, the byte unit is assumed to be K (Kilobytes).
-**Related environment variable:** ``KMP_STACKSIZE`` overrides ``GOMP_STACKSIZE``, which
-overrides ``OMP_STACKSIZE``.
+**Related environment variable:** `KMP_STACKSIZE` overrides `GOMP_STACKSIZE`, which
+overrides `OMP_STACKSIZE`.
**Default:**
-* 32-bit architectures: ``2M``
-* 64-bit architectures: ``4M``
+- 32-bit architectures: `2M`
+- 64-bit architectures: `4M`
-KMP_TOPOLOGY_METHOD
-"""""""""""""""""""
+#### KMP_TOPOLOGY_METHOD
Forces OpenMP to use a particular machine topology modeling method.
Possible values are:
-* ``all`` - Let OpenMP choose which topology method is most appropriate
+- `all` - Let OpenMP choose which topology method is most appropriate
based on the platform and possibly other environment variable settings.
-* ``cpuid_leaf31`` (x86 only) - Decodes the APIC identifiers as specified by leaf 31 of the
+- `cpuid_leaf31` (x86 only) - Decodes the APIC identifiers as specified by leaf 31 of the
cpuid instruction. The runtime will produce an error if the machine does not support leaf 31.
-* ``cpuid_leaf11`` (x86 only) - Decodes the APIC identifiers as specified by leaf 11 of the
+- `cpuid_leaf11` (x86 only) - Decodes the APIC identifiers as specified by leaf 11 of the
cpuid instruction. The runtime will produce an error if the machine does not support leaf 11.
-* ``cpuid_leaf4`` (x86 only) - Decodes the APIC identifiers as specified in leaf 4
+- `cpuid_leaf4` (x86 only) - Decodes the APIC identifiers as specified in leaf 4
of the cpuid instruction. The runtime will produce an error if the machine does not support leaf 4.
-* ``cpuinfo`` - If ``KMP_CPUINFO_FILE`` is not specified, forces OpenMP to
- parse :file:`/proc/cpuinfo` to determine the topology (Linux only).
- If ``KMP_CPUINFO_FILE`` is specified as described above, uses it (Windows or Linux).
-* ``group`` - Models the machine as a 2-level map, with level 0 specifying the
+- `cpuinfo` - If `KMP_CPUINFO_FILE` is not specified, forces OpenMP to
+ parse {file}`/proc/cpuinfo` to determine the topology (Linux only).
+ If `KMP_CPUINFO_FILE` is specified as described above, uses it (Windows or Linux).
+- `group` - Models the machine as a 2-level map, with level 0 specifying the
different processors in a group, and level 1 specifying the different
groups (Windows 64-bit only).
-.. note::
- Support for group is now deprecated and will be removed in a future release. Use all instead.
+:::{note}
+Support for group is now deprecated and will be removed in a future release. Use all instead.
+:::
-* ``flat`` - Models the machine as a flat (linear) list of processors.
-* ``hwloc`` - Models the machine as the Portable Hardware Locality (hwloc) library does.
+- `flat` - Models the machine as a flat (linear) list of processors.
+- `hwloc` - Models the machine as the Portable Hardware Locality (hwloc) library does.
This model is the most detailed and includes, but is not limited to: numa domains,
packages, cores, hardware threads, caches, and Windows processor groups. This method is
only available if you have configured libomp to use hwloc during CMake configuration.
**Default:** all
-KMP_VERSION
-"""""""""""
+#### KMP_VERSION
-Enables (``true``) or disables (``false``) the printing of OpenMP run-time
+Enables (`true`) or disables (`false`) the printing of OpenMP run-time
library version information during program execution.
-**Default:** ``false``
+**Default:** `false`
-KMP_WARNINGS
-""""""""""""
+#### KMP_WARNINGS
-Enables (``true``) or disables (``false``) displaying warnings from the
+Enables (`true`) or disables (`false`) displaying warnings from the
OpenMP run-time library during program execution.
-**Default:** ``true``
+**Default:** `true`
-.. _libomptarget:
+(libomptarget)=
-LLVM/OpenMP Target Host Runtime (``libomptarget``)
---------------------------------------------------
+## LLVM/OpenMP Target Host Runtime (`libomptarget`)
-.. _libopenmptarget_environment_vars:
+(libopenmptarget-environment-vars)=
-Environment Variables
-^^^^^^^^^^^^^^^^^^^^^
+### Environment Variables
-``libomptarget`` uses environment variables to control different features of the
+`libomptarget` uses environment variables to control different features of the
library at runtime. This allows the user to obtain useful runtime information as
well as enable or disable certain features. A full list of supported environment
variables is defined below.
- * ``LIBOMPTARGET_DEBUG=<Num>``
- * ``LIBOMPTARGET_PROFILE=<Filename>``
- * ``LIBOMPTARGET_PROFILE_GRANULARITY=<Num> (default 500, in us)``
- * ``LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD=<Num>``
- * ``LIBOMPTARGET_INFO=<Num>``
- * ``LIBOMPTARGET_HEAP_SIZE=<Num>``
- * ``LIBOMPTARGET_STACK_SIZE=<Num>``
- * ``LIBOMPTARGET_MAP_FORCE_ATOMIC=[TRUE/FALSE] (default TRUE)``
- * ``LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=[TRUE/FALSE] (default FALSE)``
- * ``LIBOMPTARGET_JIT_OPT_LEVEL={0,1,2,3} (default 3)``
- * ``LIBOMPTARGET_JIT_SKIP_OPT=[TRUE/FALSE] (default FALSE)``
- * ``LIBOMPTARGET_JIT_REPLACEMENT_OBJECT=<in:Filename> (object file)``
- * ``LIBOMPTARGET_JIT_REPLACEMENT_MODULE=<in:Filename> (LLVM-IR file)``
- * ``LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE=<out:Filename> (LLVM-IR file)``
- * ``LIBOMPTARGET_JIT_POST_OPT_IR_MODULE=<out:Filename> (LLVM-IR file)``
- * ``LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME=<out:Filename> (device image file)``
- * ``LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT=<Num> (default: 32)``
- * ``LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT=[TRUE/FALSE] (default TRUE)``
- * ``OFFLOAD_TRACK_ALLOCATION_TRACES=[TRUE/FALSE] (default FALSE)``
- * ``OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=<Num> (default 0)``
-
-LIBOMPTARGET_DEBUG
-""""""""""""""""""
-
-``LIBOMPTARGET_DEBUG`` controls whether or not debugging information will be
-displayed. This feature is only available if ``libomptarget`` was built with
-``-DOMPTARGET_DEBUG``. The debugging output provided is intended for use by
-``libomptarget`` developers. More user-friendly output is presented when using
-``LIBOMPTARGET_INFO``.
-
-LIBOMPTARGET_PROFILE
-""""""""""""""""""""
-
-``LIBOMPTARGET_PROFILE`` allows ``libomptarget`` to generate time profile output
-similar to Clang's ``-ftime-trace`` option. This generates a JSON file based on
-`Chrome Tracing`_ that can be viewed with ``chrome://tracing`` or the
-`Speedscope App`_. The output will be saved to the filename specified by the
-environment variable. For multi-threaded applications, profiling in ``libomp``
-is also needed. Setting the CMake option ``OPENMP_ENABLE_LIBOMP_PROFILING=ON``
-to enable the feature. This feature depends on the `LLVM Support Library`_
-for time trace output. Note that this will turn ``libomp`` into a C++ library.
-
-.. _`Chrome Tracing`: https://www.chromium.org/developers/how-tos/trace-event-profiling-tool
-
-.. _`Speedscope App`: https://www.speedscope.app/
-
-.. _`LLVM Support Library`: https://llvm.org/docs/SupportLibrary.html
-
-LIBOMPTARGET_PROFILE_GRANULARITY
-""""""""""""""""""""""""""""""""
-
-``LIBOMPTARGET_PROFILE_GRANULARITY`` allows to change the time profile
+> - `LIBOMPTARGET_DEBUG=<Num>`
+> - `LIBOMPTARGET_PROFILE=<Filename>`
+> - `LIBOMPTARGET_PROFILE_GRANULARITY=<Num> (default 500, in us)`
+> - `LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD=<Num>`
+> - `LIBOMPTARGET_INFO=<Num>`
+> - `LIBOMPTARGET_HEAP_SIZE=<Num>`
+> - `LIBOMPTARGET_STACK_SIZE=<Num>`
+> - `LIBOMPTARGET_MAP_FORCE_ATOMIC=[TRUE/FALSE] (default TRUE)`
+> - `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS=[TRUE/FALSE] (default FALSE)`
+> - `LIBOMPTARGET_JIT_OPT_LEVEL={0,1,2,3} (default 3)`
+> - `LIBOMPTARGET_JIT_SKIP_OPT=[TRUE/FALSE] (default FALSE)`
+> - `LIBOMPTARGET_JIT_REPLACEMENT_OBJECT=<in:Filename> (object file)`
+> - `LIBOMPTARGET_JIT_REPLACEMENT_MODULE=<in:Filename> (LLVM-IR file)`
+> - `LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE=<out:Filename> (LLVM-IR file)`
+> - `LIBOMPTARGET_JIT_POST_OPT_IR_MODULE=<out:Filename> (LLVM-IR file)`
+> - `LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME=<out:Filename> (device image file)`
+> - `LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT=<Num> (default: 32)`
+> - `LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT=[TRUE/FALSE] (default TRUE)`
+> - `OFFLOAD_TRACK_ALLOCATION_TRACES=[TRUE/FALSE] (default FALSE)`
+> - `OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=<Num> (default 0)`
+
+#### LIBOMPTARGET_DEBUG
+
+`LIBOMPTARGET_DEBUG` controls whether or not debugging information will be
+displayed. This feature is only available if `libomptarget` was built with
+`-DOMPTARGET_DEBUG`. The debugging output provided is intended for use by
+`libomptarget` developers. More user-friendly output is presented when using
+`LIBOMPTARGET_INFO`.
+
+#### LIBOMPTARGET_PROFILE
+
+`LIBOMPTARGET_PROFILE` allows `libomptarget` to generate time profile output
+similar to Clang's `-ftime-trace` option. This generates a JSON file based on
+[Chrome Tracing][chrome tracing] that can be viewed with `chrome://tracing` or the
+[Speedscope App][speedscope app]. The output will be saved to the filename specified by the
+environment variable. For multi-threaded applications, profiling in `libomp`
+is also needed. Setting the CMake option `OPENMP_ENABLE_LIBOMP_PROFILING=ON`
+to enable the feature. This feature depends on the [LLVM Support Library][llvm support library]
+for time trace output. Note that this will turn `libomp` into a C++ library.
+
+#### LIBOMPTARGET_PROFILE_GRANULARITY
+
+`LIBOMPTARGET_PROFILE_GRANULARITY` allows to change the time profile
granularity measured in `us`. Default is 500 (`us`).
-LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD
-"""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD
-``LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD`` sets the threshold size for which the
-``libomptarget`` device memory manager will handle the allocation. Any
+`LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD` sets the threshold size for which the
+`libomptarget` device 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 device memory
+freed after the device kernel exits. The default threshold value is `8KB`. If
+`LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD` is set to `0` the device memory
manager will be completely disabled.
This has no effect on the host or shared memory managers.
-.. _libomptarget_info:
+(libomptarget-info)=
-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
+`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
@@ -805,164 +926,163 @@ variable declarations in the information messages. OpenMP Debugging information
is enabled at any level of debugging so a full debug runtime is not required.
For minimal debugging information compile with `-gline-tables-only`, or compile
with `-g` for full debug information. A full list of flags supported by
-``LIBOMPTARGET_INFO`` is given below.
+`LIBOMPTARGET_INFO` is given below.
- * Print all data arguments upon entering an OpenMP device kernel: ``0x01``
- * Indicate when a mapped address already exists in the device mapping table:
- ``0x02``
- * Dump the contents of the device pointer map at kernel exit: ``0x04``
- * Indicate when an entry is changed in the device mapping table: ``0x08``
- * Print OpenMP kernel information from device plugins: ``0x10``
- * Indicate when data is copied to and from the device: ``0x20``
+> - Print all data arguments upon entering an OpenMP device kernel: `0x01`
+> - Indicate when a mapped address already exists in the device mapping table:
+> `0x02`
+> - Dump the contents of the device pointer map at kernel exit: `0x04`
+> - Indicate when an entry is changed in the device mapping table: `0x08`
+> - Print OpenMP kernel information from device plugins: `0x10`
+> - Indicate when data is copied to and from the device: `0x20`
Any combination of these flags can be used by setting the appropriate bits. For
example, to enable printing all data active in an OpenMP target region along
-with ``CUDA`` information, run the following ``bash`` command.
-
-.. code-block:: console
+with `CUDA` information, run the following `bash` command.
- $ env LIBOMPTARGET_INFO=$((0x1 | 0x10)) ./your-application
+```console
+$ env LIBOMPTARGET_INFO=$((0x1 | 0x10)) ./your-application
+```
Or, to enable every flag run with every bit set.
-.. code-block:: console
+```console
+$ env LIBOMPTARGET_INFO=-1 ./your-application
+```
- $ env LIBOMPTARGET_INFO=-1 ./your-application
-
-For example, given a small application implementing the ``ZAXPY`` BLAS routine,
-``Libomptarget`` can provide useful information about data mappings and thread
+For example, given a small application implementing the `ZAXPY` BLAS routine,
+`Libomptarget` can provide useful information about data mappings and thread
usages.
-.. code-block:: c++
-
- #include <complex>
+```c++
+#include <complex>
- using complex = std::complex<double>;
+using complex = std::complex<double>;
- void zaxpy(complex *X, complex *Y, complex D, std::size_t N) {
- #pragma omp target teams distribute parallel for
- for (std::size_t i = 0; i < N; ++i)
- Y[i] = D * X[i] + Y[i];
- }
+void zaxpy(complex *X, complex *Y, complex D, std::size_t N) {
+#pragma omp target teams distribute parallel for
+ for (std::size_t i = 0; i < N; ++i)
+ Y[i] = D * X[i] + Y[i];
+}
- int main() {
- const std::size_t N = 1024;
- complex X[N], Y[N], D;
- #pragma omp target data map(to:X[0 : N]) map(tofrom:Y[0 : N])
- zaxpy(X, Y, D, N);
- }
+int main() {
+ const std::size_t N = 1024;
+ complex X[N], Y[N], D;
+#pragma omp target data map(to:X[0 : N]) map(tofrom:Y[0 : N])
+ zaxpy(X, Y, D, N);
+}
+```
-Compiling this code targeting ``nvptx64`` with all information enabled will
+Compiling this code targeting `nvptx64` with all information enabled will
provide the following output from the runtime library.
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 -gline-tables-only zaxpy.cpp -o zaxpy
- $ env LIBOMPTARGET_INFO=-1 ./zaxpy
-
-.. code-block:: text
-
- Info: Entering OpenMP data region at zaxpy.cpp:14:1 with 2 arguments:
- Info: to(X[0:N])[16384]
- Info: tofrom(Y[0:N])[16384]
- Info: Creating new map entry with HstPtrBegin=0x00007fff0d259a40,
- TgtPtrBegin=0x00007fdba5800000, Size=16384, RefCount=1, Name=X[0:N]
- Info: Copying data from host to device, HstPtr=0x00007fff0d259a40,
- TgtPtr=0x00007fdba5800000, Size=16384, Name=X[0:N]
- Info: Creating new map entry with HstPtrBegin=0x00007fff0d255a40,
- TgtPtrBegin=0x00007fdba5804000, Size=16384, RefCount=1, Name=Y[0:N]
- Info: Copying data from host to device, HstPtr=0x00007fff0d255a40,
- TgtPtr=0x00007fdba5804000, Size=16384, Name=Y[0:N]
- Info: OpenMP Host-Device pointer mappings after block at zaxpy.cpp:14:1:
- Info: Host Ptr Target Ptr Size (B) RefCount Declaration
- Info: 0x00007fff0d255a40 0x00007fdba5804000 16384 1 Y[0:N] at zaxpy.cpp:13:17
- Info: 0x00007fff0d259a40 0x00007fdba5800000 16384 1 X[0:N] at zaxpy.cpp:13:11
- Info: Entering OpenMP kernel at zaxpy.cpp:6:1 with 4 arguments:
- Info: firstprivate(N)[8] (implicit)
- Info: use_address(Y)[0] (implicit)
- Info: tofrom(D)[16] (implicit)
- Info: use_address(X)[0] (implicit)
- Info: Mapping exists (implicit) with HstPtrBegin=0x00007fff0d255a40,
- TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=2 (incremented), Name=Y
- Info: Creating new map entry with HstPtrBegin=0x00007fff0d2559f0,
- TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1, Name=D
- Info: Copying data from host to device, HstPtr=0x00007fff0d2559f0,
- TgtPtr=0x00007fdba5808000, Size=16, Name=D
- Info: Mapping exists (implicit) with HstPtrBegin=0x00007fff0d259a40,
- TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=2 (incremented), Name=X
- Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
- TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=2 (update suppressed)
- Info: Mapping exists with HstPtrBegin=0x00007fff0d2559f0,
- TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1 (update suppressed)
- Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
- TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=2 (update suppressed)
- Info: Launching kernel __omp_offloading_10305_c08c86__Z5zaxpyPSt7complexIdES1_S0_m_l6
- with 8 blocks and 128 threads in SPMD mode
- Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
- TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=1 (decremented)
- Info: Mapping exists with HstPtrBegin=0x00007fff0d2559f0,
- TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1 (deferred final decrement)
- Info: Copying data from device to host, TgtPtr=0x00007fdba5808000,
- HstPtr=0x00007fff0d2559f0, Size=16, Name=D
- Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
- TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=1 (decremented)
- Info: Removing map entry with HstPtrBegin=0x00007fff0d2559f0,
- TgtPtrBegin=0x00007fdba5808000, Size=16, Name=D
- Info: OpenMP Host-Device pointer mappings after block at zaxpy.cpp:6:1:
- Info: Host Ptr Target Ptr Size (B) RefCount Declaration
- Info: 0x00007fff0d255a40 0x00007fdba5804000 16384 1 Y[0:N] at zaxpy.cpp:13:17
- Info: 0x00007fff0d259a40 0x00007fdba5800000 16384 1 X[0:N] at zaxpy.cpp:13:11
- Info: Exiting OpenMP data region at zaxpy.cpp:14:1 with 2 arguments:
- Info: to(X[0:N])[16384]
- Info: tofrom(Y[0:N])[16384]
- Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
- TgtPtrBegin=0x00007fdba5804000, Size=16384, RefCount=1 (deferred final decrement)
- Info: Copying data from device to host, TgtPtr=0x00007fdba5804000,
- HstPtr=0x00007fff0d255a40, Size=16384, Name=Y[0:N]
- Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
- TgtPtrBegin=0x00007fdba5800000, Size=16384, RefCount=1 (deferred final decrement)
- Info: Removing map entry with HstPtrBegin=0x00007fff0d255a40,
- TgtPtrBegin=0x00007fdba5804000, Size=16384, Name=Y[0:N]
- Info: Removing map entry with HstPtrBegin=0x00007fff0d259a40,
- TgtPtrBegin=0x00007fdba5800000, Size=16384, Name=X[0:N]
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 -gline-tables-only zaxpy.cpp -o zaxpy
+$ env LIBOMPTARGET_INFO=-1 ./zaxpy
+```
+
+```text
+Info: Entering OpenMP data region at zaxpy.cpp:14:1 with 2 arguments:
+Info: to(X[0:N])[16384]
+Info: tofrom(Y[0:N])[16384]
+Info: Creating new map entry with HstPtrBegin=0x00007fff0d259a40,
+ TgtPtrBegin=0x00007fdba5800000, Size=16384, RefCount=1, Name=X[0:N]
+Info: Copying data from host to device, HstPtr=0x00007fff0d259a40,
+ TgtPtr=0x00007fdba5800000, Size=16384, Name=X[0:N]
+Info: Creating new map entry with HstPtrBegin=0x00007fff0d255a40,
+ TgtPtrBegin=0x00007fdba5804000, Size=16384, RefCount=1, Name=Y[0:N]
+Info: Copying data from host to device, HstPtr=0x00007fff0d255a40,
+ TgtPtr=0x00007fdba5804000, Size=16384, Name=Y[0:N]
+Info: OpenMP Host-Device pointer mappings after block at zaxpy.cpp:14:1:
+Info: Host Ptr Target Ptr Size (B) RefCount Declaration
+Info: 0x00007fff0d255a40 0x00007fdba5804000 16384 1 Y[0:N] at zaxpy.cpp:13:17
+Info: 0x00007fff0d259a40 0x00007fdba5800000 16384 1 X[0:N] at zaxpy.cpp:13:11
+Info: Entering OpenMP kernel at zaxpy.cpp:6:1 with 4 arguments:
+Info: firstprivate(N)[8] (implicit)
+Info: use_address(Y)[0] (implicit)
+Info: tofrom(D)[16] (implicit)
+Info: use_address(X)[0] (implicit)
+Info: Mapping exists (implicit) with HstPtrBegin=0x00007fff0d255a40,
+ TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=2 (incremented), Name=Y
+Info: Creating new map entry with HstPtrBegin=0x00007fff0d2559f0,
+ TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1, Name=D
+Info: Copying data from host to device, HstPtr=0x00007fff0d2559f0,
+ TgtPtr=0x00007fdba5808000, Size=16, Name=D
+Info: Mapping exists (implicit) with HstPtrBegin=0x00007fff0d259a40,
+ TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=2 (incremented), Name=X
+Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
+ TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=2 (update suppressed)
+Info: Mapping exists with HstPtrBegin=0x00007fff0d2559f0,
+ TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1 (update suppressed)
+Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
+ TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=2 (update suppressed)
+Info: Launching kernel __omp_offloading_10305_c08c86__Z5zaxpyPSt7complexIdES1_S0_m_l6
+ with 8 blocks and 128 threads in SPMD mode
+Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
+ TgtPtrBegin=0x00007fdba5800000, Size=0, RefCount=1 (decremented)
+Info: Mapping exists with HstPtrBegin=0x00007fff0d2559f0,
+ TgtPtrBegin=0x00007fdba5808000, Size=16, RefCount=1 (deferred final decrement)
+Info: Copying data from device to host, TgtPtr=0x00007fdba5808000,
+ HstPtr=0x00007fff0d2559f0, Size=16, Name=D
+Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
+ TgtPtrBegin=0x00007fdba5804000, Size=0, RefCount=1 (decremented)
+Info: Removing map entry with HstPtrBegin=0x00007fff0d2559f0,
+ TgtPtrBegin=0x00007fdba5808000, Size=16, Name=D
+Info: OpenMP Host-Device pointer mappings after block at zaxpy.cpp:6:1:
+Info: Host Ptr Target Ptr Size (B) RefCount Declaration
+Info: 0x00007fff0d255a40 0x00007fdba5804000 16384 1 Y[0:N] at zaxpy.cpp:13:17
+Info: 0x00007fff0d259a40 0x00007fdba5800000 16384 1 X[0:N] at zaxpy.cpp:13:11
+Info: Exiting OpenMP data region at zaxpy.cpp:14:1 with 2 arguments:
+Info: to(X[0:N])[16384]
+Info: tofrom(Y[0:N])[16384]
+Info: Mapping exists with HstPtrBegin=0x00007fff0d255a40,
+ TgtPtrBegin=0x00007fdba5804000, Size=16384, RefCount=1 (deferred final decrement)
+Info: Copying data from device to host, TgtPtr=0x00007fdba5804000,
+ HstPtr=0x00007fff0d255a40, Size=16384, Name=Y[0:N]
+Info: Mapping exists with HstPtrBegin=0x00007fff0d259a40,
+ TgtPtrBegin=0x00007fdba5800000, Size=16384, RefCount=1 (deferred final decrement)
+Info: Removing map entry with HstPtrBegin=0x00007fff0d255a40,
+ TgtPtrBegin=0x00007fdba5804000, Size=16384, Name=Y[0:N]
+Info: Removing map entry with HstPtrBegin=0x00007fff0d259a40,
+ TgtPtrBegin=0x00007fdba5800000, Size=16384, Name=X[0:N]
+```
From this information, we can see the OpenMP kernel being launched on the CUDA
-device with enough threads and blocks for all ``1024`` iterations of the loop in
-simplified :doc:`SPMD Mode <Offloading>`. The information from the OpenMP data
-region shows the two arrays ``X`` and ``Y`` being copied from the host to the
+device with enough threads and blocks for all `1024` iterations of the loop in
+simplified {doc}`SPMD Mode <Offloading>`. The information from the OpenMP data
+region shows the two arrays `X` and `Y` being copied from the host to the
device. This creates an entry in the host-device mapping table associating the
host pointers to the newly created device data. The data mappings in the OpenMP
device kernel show the default mappings being used for all the variables used
-implicitly on the device. Because ``X`` and ``Y`` are already mapped in the
+implicitly on the device. Because `X` and `Y` are already mapped in the
device's table, no new entries are created. Additionally, the default mapping
-shows that ``D`` will be copied back from the device once the OpenMP device
+shows that `D` will be copied back from the device once the OpenMP device
kernel region ends even though it isn't written to. Finally, at the end of the
-OpenMP data region the entries for ``X`` and ``Y`` are removed from the table.
+OpenMP data region the entries for `X` and `Y` are removed from the table.
The information level can be controlled at runtime using an internal
-libomptarget library call ``__tgt_set_info_flag``. This allows for different
+libomptarget library call `__tgt_set_info_flag`. This allows for different
levels of information to be enabled or disabled for certain regions of code.
Using this requires declaring the function signature as an external function so
it can be linked with the runtime library.
-.. code-block:: c++
+```c++
+extern "C" void __tgt_set_info_flag(uint32_t);
- extern "C" void __tgt_set_info_flag(uint32_t);
+extern foo();
- extern foo();
+int main() {
+ __tgt_set_info_flag(0x10);
+#pragma omp target
+ foo();
+}
+```
- int main() {
- __tgt_set_info_flag(0x10);
- #pragma omp target
- foo();
- }
+(libopenmptarget-errors)=
-.. _libopenmptarget_errors:
+### Errors:
-Errors:
-^^^^^^^
-
-``libomptarget`` provides error messages when the program fails inside the
+`libomptarget` provides error messages when the program fails inside the
OpenMP target region. Common causes of failure could be an invalid pointer
access, running out of device memory, or trying to offload when the device is
busy. If the application was built with debugging symbols the error messages
@@ -972,103 +1092,100 @@ For example, consider the following code that implements a simple parallel
reduction on the GPU. This code has a bug that causes it to fail in the
offloading region.
-.. code-block:: c++
-
- #include <cstdio>
+```c++
+#include <cstdio>
- double sum(double *A, std::size_t N) {
- double sum = 0.0;
- #pragma omp target teams distribute parallel for reduction(+:sum)
- for (int i = 0; i < N; ++i)
- sum += A[i];
+double sum(double *A, std::size_t N) {
+ double sum = 0.0;
+#pragma omp target teams distribute parallel for reduction(+:sum)
+ for (int i = 0; i < N; ++i)
+ sum += A[i];
- return sum;
- }
+ return sum;
+}
- int main() {
- const int N = 1024;
- double A[N];
- sum(A, N);
- }
+int main() {
+ const int N = 1024;
+ double A[N];
+ sum(A, N);
+}
+```
If this code is compiled and run, there will be an error message indicating what is
going wrong.
-.. code-block:: console
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 -gline-tables-only sum.cpp -o sum
+$ ./sum
+```
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 -gline-tables-only sum.cpp -o sum
- $ ./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: 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
+```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: 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 occurring 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
+`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
time of failure.
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 -gline-tables-only sum.cpp -o sum
- $ env LIBOMPTARGET_INFO=4 ./sum
-
-.. code-block:: text
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 -gline-tables-only sum.cpp -o sum
+$ env LIBOMPTARGET_INFO=4 ./sum
+```
- info: OpenMP Host-Device pointer mappings after block at sum.cpp:5:1:
- info: Host Ptr Target Ptr Size (B) RefCount Declaration
- info: 0x00007ffc058280f8 0x00007f4186600000 8 1 sum at sum.cpp:4:10
+```text
+info: OpenMP Host-Device pointer mappings after block at sum.cpp:5:1:
+info: Host Ptr Target Ptr Size (B) RefCount Declaration
+info: 0x00007ffc058280f8 0x00007f4186600000 8 1 sum at sum.cpp:4:10
+```
This tells us that the only data mapped between the host and the device is the
-``sum`` variable that will be copied back from the device once the reduction has
-ended. There is no entry mapping the host array ``A`` to the device. In this
+`sum` variable that will be copied back from the device once the reduction has
+ended. There is no entry mapping the host array `A` to the device. In this
situation, the compiler cannot determine the size of the array at compile time
so it will simply assume that the pointer is mapped on the device already by
default. The solution is to add an explicit map clause in the target region.
-.. code-block:: c++
+```c++
+double sum(double *A, std::size_t N) {
+ double sum = 0.0;
+#pragma omp target teams distribute parallel for reduction(+:sum) map(to:A[0 : N])
+ for (int i = 0; i < N; ++i)
+ sum += A[i];
- double sum(double *A, std::size_t N) {
- double sum = 0.0;
- #pragma omp target teams distribute parallel for reduction(+:sum) map(to:A[0 : N])
- for (int i = 0; i < N; ++i)
- sum += A[i];
+ return sum;
+}
+```
- return sum;
- }
-
-LIBOMPTARGET_STACK_SIZE
-"""""""""""""""""""""""
+#### LIBOMPTARGET_STACK_SIZE
This environment variable sets the stack size in bytes for the AMDGPU and CUDA
plugins. This can be used to increase or decrease the standard amount of memory
reserved for each thread's stack.
-LIBOMPTARGET_HEAP_SIZE
-"""""""""""""""""""""""
+#### LIBOMPTARGET_HEAP_SIZE
This environment variable sets the amount of memory in bytes that can be
-allocated using ``malloc`` and ``free`` for the CUDA plugin. This is necessary
+allocated using `malloc` and `free` for the CUDA plugin. This is necessary
for some applications that allocate too much memory either through the user or
globalization.
-.. toctree::
- :hidden:
- :maxdepth: 1
-
- Offloading
+```{toctree}
+:hidden: true
+:maxdepth: 1
+Offloading
+```
-LIBOMPTARGET_MAP_FORCE_ATOMIC
-"""""""""""""""""""""""""""""
+#### LIBOMPTARGET_MAP_FORCE_ATOMIC
The OpenMP standard guarantees that map clauses are atomic. However, the this
can have a drastic performance impact. Users that do not require atomic map
@@ -1078,45 +1195,41 @@ 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.
+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_TREAT_ATTACH_AUTO_AS_ALWAYS
-"""""""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS
By default, OpenMP attach operations only perform pointer attachment
when mapping an expression with a base-pointer/base-referring-pointer,
when either the pointer or the pointee was newly allocated on a
-map-entering directive (aka ``attach(auto)`` as per OpenMP 6.1 TR14).
+map-entering directive (aka `attach(auto)` as per OpenMP 6.1 TR14).
-When ``LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS`` is set to ``true``,
+When `LIBOMPTARGET_TREAT_ATTACH_AUTO_AS_ALWAYS` is set to `true`,
ATTACH map entries without the ALWAYS flag are implicitly treated as if
the ALWAYS flag was set. This forces pointer attachments to occur even when
the pointee/pointer were not newly allocated (similar to OpenMP 6.1
-TR14's ``attach(always)`` map-type-modifier), thereby treating
-``attach(auto))`` as ``attach(always)``. This can be used for
+TR14's `attach(always)` map-type-modifier), thereby treating
+`attach(auto))` as `attach(always)`. This can be used for
experimentation, or as a workaround for programs compiled without
-``-fopenmp-version=61``.
+`-fopenmp-version=61`.
-.. _libomptarget_jit_opt_level:
+(libomptarget-jit-opt-level)=
-LIBOMPTARGET_JIT_OPT_LEVEL
-""""""""""""""""""""""""""
+#### LIBOMPTARGET_JIT_OPT_LEVEL
This environment variable can be used to change the optimization pipeline used
to optimize the embedded device code as part of the device JIT. The value is
-corresponds to the ``-O{0,1,2,3}`` command line argument passed to ``clang``.
+corresponds to the `-O{0,1,2,3}` command line argument passed to `clang`.
-LIBOMPTARGET_JIT_SKIP_OPT
-""""""""""""""""""""""""""
+#### LIBOMPTARGET_JIT_SKIP_OPT
This environment variable can be used to skip the optimization pipeline during
JIT compilation. If set, the image will only be passed through the backend. The
-backend is invoked with the ``LIBOMPTARGET_JIT_OPT_LEVEL`` flag.
+backend is invoked with the `LIBOMPTARGET_JIT_OPT_LEVEL` flag.
-LIBOMPTARGET_JIT_REPLACEMENT_OBJECT
-"""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_JIT_REPLACEMENT_OBJECT
This environment variable can be used to replace the embedded device code
before the device JIT finishes compilation for the target. The value is
@@ -1125,10 +1238,9 @@ assembler in object format for the respective target. The JIT optimization
pipeline and backend are skipped and only target specific post-processing is
performed on the object file before it is loaded onto the device.
-.. _libomptarget_jit_replacement_module:
+(libomptarget-jit-replacement-module)=
-LIBOMPTARGET_JIT_REPLACEMENT_MODULE
-"""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_JIT_REPLACEMENT_MODULE
This environment variable can be used to replace the embedded device code
before the device JIT finishes compilation for the target. The value is
@@ -1137,81 +1249,73 @@ for the respective target. To obtain a device code image compatible with the
embedded one it is recommended to extract the embedded one either before or
after IR optimization. This can be done at compile time, after compile time via
llvm tools (llvm-objdump), or, simply, by setting the
-:ref:`LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE` or
-:ref:`LIBOMPTARGET_JIT_POST_OPT_IR_MODULE` environment variables.
+{ref}`LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE` or
+{ref}`LIBOMPTARGET_JIT_POST_OPT_IR_MODULE` environment variables.
-.. _libomptarget_jit_pre_opt_ir_module:
+(libomptarget-jit-pre-opt-ir-module)=
-LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE
-""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE
This environment variable can be used to extract the embedded device code
before the device JIT runs additional IR optimizations on it (see
-:ref:`LIBOMPTARGET_JIT_OPT_LEVEL`). The value is expected to be a filename into
+{ref}`LIBOMPTARGET_JIT_OPT_LEVEL`). The value is expected to be a filename into
which the LLVM-IR module is written. The module can be the analyzed, and
transformed and loaded back into the JIT pipeline via
-:ref:`LIBOMPTARGET_JIT_REPLACEMENT_MODULE`.
+{ref}`LIBOMPTARGET_JIT_REPLACEMENT_MODULE`.
-.. _libomptarget_jit_post_opt_ir_module:
+(libomptarget-jit-post-opt-ir-module)=
-LIBOMPTARGET_JIT_POST_OPT_IR_MODULE
-"""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_JIT_POST_OPT_IR_MODULE
This environment variable can be used to extract the embedded device code after
the device JIT runs additional IR optimizations on it (see
-:ref:`LIBOMPTARGET_JIT_OPT_LEVEL`). The value is expected to be a filename into
+{ref}`LIBOMPTARGET_JIT_OPT_LEVEL`). The value is expected to be a filename into
which the LLVM-IR module is written. The module can be the analyzed, and
transformed and loaded back into the JIT pipeline via
-:ref:`LIBOMPTARGET_JIT_REPLACEMENT_MODULE`.
+{ref}`LIBOMPTARGET_JIT_REPLACEMENT_MODULE`.
-.. _libomptarget_jit_save_image_filename:
+(libomptarget-jit-save-image-filename)=
-LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME
-""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME
This environment variable can be used to save the device image produced by the
device JIT after target-specific post-processing. The value is expected to be a
filename into which the binary device image is written.
-LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT
-"""""""""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT
This environment variable defines a lower bound for the number of threads if a
combined kernel, e.g., `target teams distribute parallel for`, has insufficient
parallelism. Especially if the trip count of the loops is lower than the number
of threads possible times the number of teams (aka. blocks) the device prefers
-(see also :ref:`LIBOMPTARGET_AMDGPU_TEAMS_PER_CU`), we will reduce the thread
+(see also {ref}`LIBOMPTARGET_AMDGPU_TEAMS_PER_CU`), we will reduce the thread
count to increase outer (team/block) parallelism. The thread count will never
be reduced below the value passed for this environment variable though.
-LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT
-"""""""""""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT
This environment variable can be used to control how the OpenMP runtime assigns
blocks to loops with high trip counts. By default we reuse existing blocks
rather than spawning new blocks.
-OFFLOAD_TRACK_ALLOCATION_TRACES
-"""""""""""""""""""""""""""""""
+#### OFFLOAD_TRACK_ALLOCATION_TRACES
This environment variable determines if the stack traces of allocations and
deallocations are tracked to aid in error reporting, e.g., in case of
double-free.
-OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES
-""""""""""""""""""""""""""""""""""""""
+#### OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES
This environment variable determines how many stack traces of kernel launches
are tracked to aid in error reporting, e.g., what asynchronous kernel failed.
-.. _libomptarget_kernel_record_replay:
+(libomptarget-kernel-record-replay)=
-Kernel Record Replay
-^^^^^^^^^^^^^^^^^^^^
+### Kernel Record Replay
The Kernel Record and Replay mechanism enables recording the execution of GPU
kernels on OpenMP applications and replaying them in isolation using
-``llvm-omp-kernel-replay``, a lightweight LLVM-based tool. This tool is useful
+`llvm-omp-kernel-replay`, a lightweight LLVM-based tool. This tool is useful
for extracting kernel executions from applications to analyze them
independently, with the flexibility to modify certain runtime parameters.
@@ -1222,53 +1326,53 @@ generates a JSON file that describes the kernel alongside the runtime parameters
(e.g., the number of teams and threads).
To record the kernels of an OpenMP application, enable the
-:ref:`LIBOMPTARGET_RECORD` environment variable when running the program. An
+{ref}`LIBOMPTARGET_RECORD` environment variable when running the program. An
example is shown below:
-.. code-block:: console
-
- $ LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_REPORT=1 LIBOMPTARGET_RECORD_DIR=records ./application
- ... application output ...
- === Kernel Record Report ===
- Directory: /home/records
- Total Instances: 1
- JSON Filename, Kernel Name, Time (ns), Occurrences:
- 5681756204876336171_6652394454608725381.json, __omp_offloading_48_5f678667_run_event_based_simulation_l44, 63437836, 1
- === End Kernel Record Report ===
+```console
+$ LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_REPORT=1 LIBOMPTARGET_RECORD_DIR=records ./application
+... application output ...
+=== Kernel Record Report ===
+Directory: /home/records
+Total Instances: 1
+JSON Filename, Kernel Name, Time (ns), Occurrences:
+5681756204876336171_6652394454608725381.json, __omp_offloading_48_5f678667_run_event_based_simulation_l44, 63437836, 1
+=== End Kernel Record Report ===
+```
The command above creates a directory (as indicated by
-:ref:`LIBOMPTARGET_RECORD_DIR`) containing the memory snapshots and a JSON file
+{ref}`LIBOMPTARGET_RECORD_DIR`) containing the memory snapshots and a JSON file
for each recorded kernel. This JSON file contains the description, properties,
and original runtime parameters of the kernel. Additionally, enabling
-:ref:`LIBOMPTARGET_RECORD_REPORT` instructs the runtime to emit a summary of the
+{ref}`LIBOMPTARGET_RECORD_REPORT` instructs the runtime to emit a summary of the
recorded kernel instances and their associated JSON files.
-To replay a particular kernel, run the ``llvm-omp-kernel-replay`` command,
+To replay a particular kernel, run the `llvm-omp-kernel-replay` command,
passing the path to the corresponding kernel's JSON file:
-.. code-block:: console
-
- $ llvm-omp-kernel-replay --repetitions=5 records/5681756204876336171_6652394454608725381.json
- [llvm-omp-kernel-replay] Replay time (1): 94926702 ns
- [llvm-omp-kernel-replay] Replay time (2): 94642823 ns
- [llvm-omp-kernel-replay] Replay time (3): 94429614 ns
- [llvm-omp-kernel-replay] Replay time (4): 94574421 ns
- [llvm-omp-kernel-replay] Replay time (5): 94359425 ns
- [llvm-omp-kernel-replay] Replay done, verification skipped
+```console
+$ llvm-omp-kernel-replay --repetitions=5 records/5681756204876336171_6652394454608725381.json
+[llvm-omp-kernel-replay] Replay time (1): 94926702 ns
+[llvm-omp-kernel-replay] Replay time (2): 94642823 ns
+[llvm-omp-kernel-replay] Replay time (3): 94429614 ns
+[llvm-omp-kernel-replay] Replay time (4): 94574421 ns
+[llvm-omp-kernel-replay] Replay time (5): 94359425 ns
+[llvm-omp-kernel-replay] Replay done, verification skipped
+```
When replaying, you can tune the execution using the following flags, among
others:
-* ``--repetitions=N``: Sets the number of repetitions of the kernel replay
+- `--repetitions=N`: Sets the number of repetitions of the kernel replay
(default 1).
-* ``--num-threads=N``: Overrides the number of threads per team.
-* ``--num-teams=N``: Overrides the number of teams.
+- `--num-threads=N`: Overrides the number of threads per team.
+- `--num-teams=N`: Overrides the number of teams.
-If ``--num-threads`` or ``--num-teams`` are not specified, the replay
+If `--num-threads` or `--num-teams` are not specified, the replay
automatically defaults to the values used during the original recorded run. The
replay tool will issue an error if you specify a number of threads or teams that
is incompatible with the limits established by the original code (e.g.,
-exceeding bounds set by a ``num_teams`` or ``thread_limit`` clause).
+exceeding bounds set by a `num_teams` or `thread_limit` clause).
The time reported by the replay tool corresponds to the host-side kernel launch
and synchronization time. If highly precise kernel timing is required, it is
@@ -1276,28 +1380,27 @@ recommended to use dedicated profiling tools in conjunction with the replay
tool.
Finally, the replay tool provides an optional verification step via the
-``--verify`` flag. When enabled, it checks whether the output device memory
+`--verify` flag. When enabled, it checks whether the output device memory
snapshot generated during replay matches the output snapshot captured during the
recording phase. Because this verification performs a strict binary difference
between the two memory snapshots, the check may fail for kernels operating on
floating-point data due to normal variations in precision and operation order.
-The recording phase, implemented by ``libomptarget``, can be controlled via
+The recording phase, implemented by `libomptarget`, can be controlled via
environment variables. A full list of environment variables and their definition
is provided below.
-* ``LIBOMPTARGET_RECORD=[TRUE/FALSE] (default FALSE)``
-* ``LIBOMPTARGET_RECORD_DIR=<Filepath>``
-* ``LIBOMPTARGET_RECORD_REPORT=[TRUE/FALSE] (default FALSE)``
-* ``LIBOMPTARGET_RECORD_REPORT_FILENAME=<Filename>``
-* ``LIBOMPTARGET_RECORD_MEMSIZE=<Num> (default 8*1024*1024*1024)``
-* ``LIBOMPTARGET_RECORD_DEVICE=<Num> (default 0)``
-* ``LIBOMPTARGET_RECORD_OUTPUT=[TRUE/FALSE] (default TRUE)``
+- `LIBOMPTARGET_RECORD=[TRUE/FALSE] (default FALSE)`
+- `LIBOMPTARGET_RECORD_DIR=<Filepath>`
+- `LIBOMPTARGET_RECORD_REPORT=[TRUE/FALSE] (default FALSE)`
+- `LIBOMPTARGET_RECORD_REPORT_FILENAME=<Filename>`
+- `LIBOMPTARGET_RECORD_MEMSIZE=<Num> (default 8*1024*1024*1024)`
+- `LIBOMPTARGET_RECORD_DEVICE=<Num> (default 0)`
+- `LIBOMPTARGET_RECORD_OUTPUT=[TRUE/FALSE] (default TRUE)`
-.. _libomptarget_record:
+(libomptarget-record)=
-LIBOMPTARGET_RECORD
-"""""""""""""""""""
+#### LIBOMPTARGET_RECORD
This environment variable is used to enable the kernel recording mechanism in
the execution of a OpenMP program. Enabling the record may introduce significant
@@ -1305,62 +1408,55 @@ overhead to the recorded program. When the recording is disabled, the following
recording environment variables are not considered. The recording is disabled by
default.
-.. _libomptarget_record_dir:
+(libomptarget-record-dir)=
-LIBOMPTARGET_RECORD_DIR
-"""""""""""""""""""""""
+#### LIBOMPTARGET_RECORD_DIR
This environment variable is used to specify the relative or absolute path to
the directory where the recorded files will be stored. If omitted or empty, the
files will be stored in current working directory.
-.. _libomptarget_record_report:
+(libomptarget-record-report)=
-LIBOMPTARGET_RECORD_REPORT
-""""""""""""""""""""""""""
+#### LIBOMPTARGET_RECORD_REPORT
This environment variable is used to instruct the runtime to emit a summary of
the recorded kernel instances and their associated JSON files. When enabled, the
report is emitted in the standard output. See
-:ref:`LIBOMPTARGET_RECORD_REPORT_FILENAME` to emit the report to a file. By
+{ref}`LIBOMPTARGET_RECORD_REPORT_FILENAME` to emit the report to a file. By
default, no report is emitted.
-.. _libomptarget_record_report_filename:
+(libomptarget-record-report-filename)=
-LIBOMPTARGET_RECORD_REPORT_FILENAME
-"""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_RECORD_REPORT_FILENAME
This environment variable is used to instruct the runtime to emit the recording
report to a file with a specific file. The file is written in the recording
-directory (see :ref:`LIBOMPTARGET_RECORD_DIR`). Note that it is not needed to
-use :ref:`LIBOMPTARGET_RECORD_REPORT` when setting this environment variable.
+directory (see {ref}`LIBOMPTARGET_RECORD_DIR`). Note that it is not needed to
+use {ref}`LIBOMPTARGET_RECORD_REPORT` when setting this environment variable.
-LIBOMPTARGET_RECORD_MEMSIZE
-"""""""""""""""""""""""""""
+#### LIBOMPTARGET_RECORD_MEMSIZE
This environment variable is used to indicate the maximum size of device virtual
memory that will be captured in the snapshots during the recording phase. This
value only indicates the maximum size; the snapshot files will just contain the
actually used data. Modifying this environment variable should be needed in very
-specific cases. By default, the size is ``8*1024*1024*1024`` bytes (8 GB).
+specific cases. By default, the size is `8*1024*1024*1024` bytes (8 GB).
-LIBOMPTARGET_RECORD_DEVICE
-""""""""""""""""""""""""""
+#### LIBOMPTARGET_RECORD_DEVICE
This environment variable is used to indicate the number of the device whose
kernels should be recorded. Only the kernels executed by this device will be
-recorded. The default device is ``0``.
+recorded. The default device is `0`.
-LIBOMPTARGET_RECORD_OUTPUT
-""""""""""""""""""""""""""
+#### LIBOMPTARGET_RECORD_OUTPUT
This environment variable is used to instruct the runtime to record the output
-device memory snapshot into a file. The default value is ``TRUE``.
+device memory snapshot into a file. The default value is `TRUE`.
-.. _libomptarget_plugin:
+(libomptarget-plugin)=
-LLVM/OpenMP Target Host Runtime Plugins (``libomptarget.rtl.XXXX``)
--------------------------------------------------------------------
+## LLVM/OpenMP Target Host Runtime Plugins (`libomptarget.rtl.XXXX`)
The LLVM/OpenMP target host runtime plugins were recently re-implemented,
temporarily renamed as the NextGen plugins, and set as the default and only
@@ -1368,7 +1464,7 @@ plugins' implementation. Currently, these plugins have support for the NVIDIA
and AMDGPU devices as well as the GenericELF64bit host-simulated device.
The source code of the common infrastructure and the vendor-specific plugins is
-in the ``openmp/libomptarget/nextgen-plugins`` directory in the LLVM project
+in the `openmp/libomptarget/nextgen-plugins` directory in the LLVM project
repository. The plugin infrastructure aims at unifying the plugin code and logic
into a generic interface using object-oriented C++. There is a plugin interface
composed by multiple generic C++ classes which implement the common logic that
@@ -1382,77 +1478,70 @@ With this common plugin infrastructure, several tasks have been simplified:
adding a new vendor-specific plugin, adding generic features or optimizations
to all plugins, debugging plugins, etc.
-Environment Variables
-^^^^^^^^^^^^^^^^^^^^^
+### Environment Variables
There are several environment variables to change the behavior of the plugins:
-* ``LIBOMPTARGET_STACK_SIZE``
-* ``LIBOMPTARGET_HEAP_SIZE``
-* ``LIBOMPTARGET_NUM_INITIAL_STREAMS``
-* ``LIBOMPTARGET_NUM_INITIAL_EVENTS``
-* ``LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS``
-* ``LIBOMPTARGET_AMDGPU_NUM_HSA_QUEUES``
-* ``LIBOMPTARGET_AMDGPU_HSA_QUEUE_SIZE``
-* ``LIBOMPTARGET_AMDGPU_HSA_QUEUE_BUSY_TRACKING``
-* ``LIBOMPTARGET_AMDGPU_TEAMS_PER_CU``
-* ``LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES``
-* ``LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS``
-* ``LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT``
-
-The environment variables ``LIBOMPTARGET_STACK_SIZE`` and
-``LIBOMPTARGET_HEAP_SIZE`` are described in
-:ref:`libopenmptarget_environment_vars`.
-
-LIBOMPTARGET_NUM_INITIAL_STREAMS
-""""""""""""""""""""""""""""""""
+- `LIBOMPTARGET_STACK_SIZE`
+- `LIBOMPTARGET_HEAP_SIZE`
+- `LIBOMPTARGET_NUM_INITIAL_STREAMS`
+- `LIBOMPTARGET_NUM_INITIAL_EVENTS`
+- `LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS`
+- `LIBOMPTARGET_AMDGPU_NUM_HSA_QUEUES`
+- `LIBOMPTARGET_AMDGPU_HSA_QUEUE_SIZE`
+- `LIBOMPTARGET_AMDGPU_HSA_QUEUE_BUSY_TRACKING`
+- `LIBOMPTARGET_AMDGPU_TEAMS_PER_CU`
+- `LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES`
+- `LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS`
+- `LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT`
+
+The environment variables `LIBOMPTARGET_STACK_SIZE` and
+`LIBOMPTARGET_HEAP_SIZE` are described in
+{ref}`libopenmptarget_environment_vars`.
+
+#### LIBOMPTARGET_NUM_INITIAL_STREAMS
This environment variable sets the number of pre-created streams in the plugin
(if supported) at initialization. More streams will be created dynamically
throughout the execution if needed. A stream is a queue of asynchronous
operations (e.g., kernel launches and memory copies) that are executed
sequentially. Parallelism is achieved by featuring multiple streams. The
-``libomptarget`` leverages streams to exploit parallelism between plugin
-operations. The default value is ``1``, more streams are created as needed.
+`libomptarget` leverages streams to exploit parallelism between plugin
+operations. The default value is `1`, more streams are created as needed.
-LIBOMPTARGET_NUM_INITIAL_EVENTS
-"""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_NUM_INITIAL_EVENTS
This environment variable sets the number of pre-created events in the
plugin (if supported) at initialization. More events will be created
dynamically throughout the execution if needed. An event is used to synchronize
-a stream with another efficiently. The default value is ``1``, more events are
+a stream with another efficiently. The default value is `1`, more events are
created as needed.
-LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS
-"""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS
This environment variable indicates whether the host buffers mapped by the user
should be automatically locked/pinned by the plugin. Pinned host buffers allow
true asynchronous copies between the host and devices. Enabling this feature can
increase the performance of applications that are intensive in host-device
-memory transfers. The default value is ``false``.
+memory transfers. The default value is `false`.
-LIBOMPTARGET_AMDGPU_NUM_HSA_QUEUES
-""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_AMDGPU_NUM_HSA_QUEUES
This environment variable controls the number of HSA queues per device in the
AMDGPU plugin. An HSA queue is a runtime-allocated resource that contains an
AQL (Architected Queuing Language) packet buffer and is associated with an AQL
packet processor. HSA queues are used for inserting kernel packets to launching
kernel executions. A high number of HSA queues may degrade the performance. The
-default value is ``4``.
+default value is `4`.
-LIBOMPTARGET_AMDGPU_HSA_QUEUE_SIZE
-""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_AMDGPU_HSA_QUEUE_SIZE
This environment variable controls the size of each HSA queue in the AMDGPU
plugin. The size is the number of AQL packets an HSA queue is expected to hold.
It is also the number of AQL packets that can be pushed into each queue without
-waiting the driver to process them. The default value is ``512``.
+waiting the driver to process them. The default value is `512`.
-LIBOMPTARGET_AMDGPU_HSA_QUEUE_BUSY_TRACKING
-"""""""""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_AMDGPU_HSA_QUEUE_BUSY_TRACKING
This environment variable controls if idle HSA queues will be preferentially
assigned to streams, for example when they are requested for a kernel launch.
@@ -1460,49 +1549,44 @@ Should all queues be considered busy, a new queue is initialized and returned,
until we reach the set maximum. Otherwise, we will select the least utilized
queue. If this is disabled, each time a stream is requested a new HSA queue
will be initialized, regardless of their utilization. Additionally, queues will
-be selected using round robin selection. The default value is ``true``.
+be selected using round robin selection. The default value is `true`.
-.. _libomptarget_amdgpu_teams_per_cu:
+(libomptarget-amdgpu-teams-per-cu)=
-LIBOMPTARGET_AMDGPU_TEAMS_PER_CU
-""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_AMDGPU_TEAMS_PER_CU
This environment variable controls the default number of teams relative to the
number of compute units (CUs) of the AMDGPU device. The default number of teams
-is ``#default_teams = #teams_per_CU * #CUs``. The default value of teams per CU
-is ``4``.
+is `#default_teams = #teams_per_CU * #CUs`. The default value of teams per CU
+is `4`.
-LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES
-""""""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES
This environment variable specifies the maximum size in bytes where the memory
copies are asynchronous operations in the AMDGPU plugin. Up to this transfer
size, the memory copies are asynchronous operations pushed to the corresponding
stream. For larger transfers, they are synchronous transfers. Memory copies
involving already locked/pinned host buffers are always asynchronous. The default
-value is ``1*1024*1024`` bytes (1 MB).
+value is `1*1024*1024` bytes (1 MB).
-LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS
-"""""""""""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS
This environment variable controls the initial number of HSA signals per device
in the AMDGPU plugin. There is one resource manager of signals per device
managing several pre-created signals. These signals are mainly used by AMDGPU
streams. More HSA signals will be created dynamically throughout the execution
-if needed. The default value is ``64``.
+if needed. The default value is `64`.
-LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT
-"""""""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_AMDGPU_STREAM_BUSYWAIT
This environment variable controls the timeout hint in microseconds for the
HSA wait state within the AMDGPU plugin. For the duration of this value
the HSA runtime may busy wait. This can reduce overall latency.
-The default value is ``2000000``.
+The default value is `2000000`.
-.. _remote_offloading_plugin:
+(remote-offloading-plugin)=
-Remote Offloading Plugin:
-^^^^^^^^^^^^^^^^^^^^^^^^^
+### Remote Offloading Plugin:
The remote offloading plugin permits the execution of OpenMP target regions
on devices in remote hosts in addition to the devices connected to the local
@@ -1513,8 +1597,8 @@ server is running on the same host, each device may be identified twice:
once through the device plugins and once through the device plugins that the
server application has access to.
-This plugin consists of ``libomptarget.rtl.rpc.so`` and
-``openmp-offloading-server`` which should be running on the (remote) host. The
+This plugin consists of `libomptarget.rtl.rpc.so` and
+`openmp-offloading-server` which should be running on the (remote) host. The
server application does not have to be running on a remote host, and can
instead be used on the same host in order to debug memory mapping during offloading.
These are implemented via gRPC/protobuf so these libraries are required to
@@ -1522,59 +1606,59 @@ build and use this plugin. The server must also have access to the necessary
target-specific plugins in order to perform the offloading.
Due to the experimental nature of this plugin, the CMake variable
-``LIBOMPTARGET_ENABLE_EXPERIMENTAL_REMOTE_PLUGIN`` must be set in order to
+`LIBOMPTARGET_ENABLE_EXPERIMENTAL_REMOTE_PLUGIN` must be set in order to
build this plugin. For example, the rpc plugin is not designed to be
thread-safe, the server cannot concurrently handle offloading from multiple
applications at once (it is synchronous) and will terminate after a single
-execution. Note that ``openmp-offloading-server`` is unable to
+execution. Note that `openmp-offloading-server` is unable to
remote offload onto a remote host itself and will error out if this is attempted.
Remote offloading is configured via environment variables at runtime of the OpenMP application:
- * ``LIBOMPTARGET_RPC_ADDRESS=<Address>:<Port>``
- * ``LIBOMPTARGET_RPC_ALLOCATOR_MAX=<NumBytes>``
- * ``LIBOMPTARGET_BLOCK_SIZE=<NumBytes>``
- * ``LIBOMPTARGET_RPC_LATENCY=<Seconds>``
+: - `LIBOMPTARGET_RPC_ADDRESS=<Address>:<Port>`
+ - `LIBOMPTARGET_RPC_ALLOCATOR_MAX=<NumBytes>`
+ - `LIBOMPTARGET_BLOCK_SIZE=<NumBytes>`
+ - `LIBOMPTARGET_RPC_LATENCY=<Seconds>`
+
+#### LIBOMPTARGET_RPC_ADDRESS
-LIBOMPTARGET_RPC_ADDRESS
-""""""""""""""""""""""""
The address and port at which the server is running. This needs to be set for
-the server and the application, the default is ``0.0.0.0:50051``. A single
+the server and the application, the default is `0.0.0.0:50051`. A single
OpenMP executable can offload onto multiple remote hosts by setting this to
comma-separated values of the addresses.
-LIBOMPTARGET_RPC_ALLOCATOR_MAX
-""""""""""""""""""""""""""""""
+#### LIBOMPTARGET_RPC_ALLOCATOR_MAX
+
After allocating this size, the protobuf allocator will clear. This can be set for both endpoints.
-LIBOMPTARGET_BLOCK_SIZE
-"""""""""""""""""""""""
+#### 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
-""""""""""""""""""""""""
+#### LIBOMPTARGET_RPC_LATENCY
+
This is the maximum amount of time the client will wait for a response from the server.
-.. warning::
- The ``LIBOMPTARGET_SHARED_MEMORY_SIZE`` environment variable is not
- supported anymore. Please use the ``dyn_groupprivate`` clause instead, as
- shown in :ref:`libomptarget_dynamic_shared`.
+:::{warning}
+The `LIBOMPTARGET_SHARED_MEMORY_SIZE` environment variable is not
+supported anymore. Please use the `dyn_groupprivate` clause instead, as
+shown in {ref}`libomptarget_dynamic_shared`.
+:::
-.. _libomptarget_libc:
+(libomptarget-libc)=
-LLVM/OpenMP support for C library routines
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+### 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
+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>`_.
+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
+`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.
@@ -1582,99 +1666,100 @@ 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++
+```c++
+#include <stdio.h>
- #include <stdio.h>
+int main() {
+ #pragma omp target
+ { fputs("Hello World!\n", stderr); }
+}
+```
- int main() {
- #pragma omp target
- { fputs("Hello World!\n", stderr); }
- }
-
-We can compile this using the ``libcgpu.a`` library to resolve the symbols.
+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
+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!
+```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:
+(libomptarget-device)=
-LLVM/OpenMP Target Device Runtime (``libomptarget-ARCH-SUBARCH.bc``)
---------------------------------------------------------------------
+## 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_dynamic_shared:
+(libomptarget-dynamic-shared)=
-Dynamic Shared Memory
-^^^^^^^^^^^^^^^^^^^^^
+### Dynamic Shared Memory
-The OpenMP implementation provides access to dynamic shared memory in ``target``
-regions through the ``dyn_groupprivate`` clause, introduced in OpenMP 6.1. This
+The OpenMP implementation provides access to dynamic shared memory in `target`
+regions through the `dyn_groupprivate` clause, introduced in OpenMP 6.1. This
is the preferred method to obtain dynamic shared memory. Please refer to
the OpenMP standard documentation for more information.
As an alternative, the target device runtime contains a pointer to the native
dynamic shared memory buffer. This pointer can be obtained using the
-``llvm_omp_target_dynamic_shared_alloc`` extension. If this function is called
+`llvm_omp_target_dynamic_shared_alloc` extension. If this function is called
from the host it will simply return a null pointer. In order to use this buffer
the kernel must be launched with an adequate amount of dynamic shared memory
-allocated. This can be done using the ``ompx_dyn_cgroup_mem(<N>)`` target
+allocated. This can be done using the `ompx_dyn_cgroup_mem(<N>)` target
directive clause. An example is given below.
-Please notice that the ``LIBOMPTARGET_SHARED_MEMORY_SIZE`` environment variable
+Please notice that the `LIBOMPTARGET_SHARED_MEMORY_SIZE` environment variable
is not supported anymore.
-.. code-block:: c++
-
- void foo(int N) {
- int x;
- #pragma omp target parallel map(from : x) ompx_dyn_cgroup_mem(N * sizeof(int))
- {
- int *buf = llvm_omp_target_dynamic_shared_alloc();
- if (omp_get_thread_num() == 0)
- buf[N - 1] = 1;
- #pragma omp barrier
- if (omp_get_thread_num() == 1)
- x = buf[N - 1];
- }
- assert(x == 1);
- }
-
-.. _libomptarget_device_allocator:
-
-Device Allocation
-^^^^^^^^^^^^^^^^^
-
-The device runtime supports basic runtime allocation via the ``omp_alloc``
-function. Currently, this allocates global memory for all default traits. Access
+```c++
+void foo(int N) {
+ int x;
+#pragma omp target parallel map(from : x) ompx_dyn_cgroup_mem(N * sizeof(int))
+ {
+ int *buf = llvm_omp_target_dynamic_shared_alloc();
+ if (omp_get_thread_num() == 0)
+ buf[N - 1] = 1;
+#pragma omp barrier
+ if (omp_get_thread_num() == 1)
+ x = buf[N - 1];
+ }
+ assert(x == 1);
+}
+```
+
+(libomptarget-device-allocator)=
+
+### Device Allocation
+
+The device runtime supports basic runtime allocation via the `omp_alloc`
+function. Currently, this allocates global memory for all default traits. Access
modifiers are currently not supported and return a null pointer.
-.. _libomptarget_device_debugging:
+(libomptarget-device-debugging)=
-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
+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
+`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 diagnosing common problems during offloading . ``0x4``
- * Dump device PGO counters (only if PGO on GPU is enabled). ``0x10``
+> - Enable debugging assertions in the device. `0x01`
+> - Enable diagnosing common problems during offloading . `0x4`
+> - Dump device PGO counters (only if PGO on GPU is enabled). `0x10`
+
+[chrome tracing]: https://www.chromium.org/developers/how-tos/trace-event-profiling-tool
+[llvm support library]: https://llvm.org/docs/SupportLibrary.html
+[speedscope app]: https://www.speedscope.app/
+
diff --git a/openmp/docs/index.md b/openmp/docs/index.md
index def98df57817c2..97650a31a36bd9 100644
--- a/openmp/docs/index.md
+++ b/openmp/docs/index.md
@@ -1,159 +1,155 @@
-.. title:: Welcome to the documentation of OpenMP in LLVM!
+```{title} Welcome to the documentation of OpenMP in LLVM!
+```
-.. note::
- This document is a work in progress and most of the expected content is not
- yet available. While you can expect changes, we always welcome feedback and
- additions. Please post on the `Discourse forums (Runtimes -
- OpenMP) <https://discourse.llvm.org/c/runtimes/openmp/35>`__..
+:::{note}
+This document is a work in progress and most of the expected content is not
+yet available. While you can expect changes, we always welcome feedback and
+additions. Please post on the [Discourse forums (Runtimes -
+OpenMP)](https://discourse.llvm.org/c/runtimes/openmp/35)..
+:::
-.. toctree::
- :hidden:
- :maxdepth: 1
+```{toctree}
+:hidden: true
+:maxdepth: 1
- LLVM/OpenMP Documentation <self>
- Building
+LLVM/OpenMP Documentation <self>
+Building
+```
-
-Getting Started
-===============
+# Getting Started
Building LLVM/OpenMP is fully documented on the
-:doc:`Building` page. For a quick start, we recommend the following template
+{doc}`Building` page. For a quick start, we recommend the following template
for building OpenMP with offloading support.
-.. code-block:: console
-
- $ git clone https://github.com/llvm/llvm-project.git
- $ cd llvm-project
- $ mkdir build
- $ cd build
- $ cmake ../llvm -G Ninja \
- -C ../offload/cmake/caches/Offload.cmake \
- -DCMAKE_BUILD_TYPE=Release \
- -DCMAKE_INSTALL_PREFIX=<PATH>
- $ ninja install # Builds all files and installs files to <PATH>/bin, <PATH>/lib, etc
+```console
+$ git clone https://github.com/llvm/llvm-project.git
+$ cd llvm-project
+$ mkdir build
+$ cd build
+$ cmake ../llvm -G Ninja \
+ -C ../offload/cmake/caches/Offload.cmake \
+ -DCMAKE_BUILD_TYPE=Release \
+ -DCMAKE_INSTALL_PREFIX=<PATH>
+$ ninja install # Builds all files and installs files to <PATH>/bin, <PATH>/lib, etc
+```
-LLVM/OpenMP Design & Overview
-=============================
+# LLVM/OpenMP Design & Overview
-OpenMP impacts various parts of the LLVM project, from the frontends (`Clang
-<https://clang.llvm.org/docs/OpenMPSupport.html>`_ and Flang), through
-middle-end :ref:`optimizations <llvm_openmp_optimizations>`, up to the
-multitude of available :ref:`OpenMP runtimes <openmp_runtimes>`.
+OpenMP impacts various parts of the LLVM project, from the frontends ([Clang](https://clang.llvm.org/docs/OpenMPSupport.html) and Flang), through
+middle-end {ref}`optimizations <llvm_openmp_optimizations>`, up to the
+multitude of available {ref}`OpenMP runtimes <openmp_runtimes>`.
-A high-level overview of OpenMP in LLVM can be found :doc:`here <design/Overview>`.
+A high-level overview of OpenMP in LLVM can be found {doc}`here <design/Overview>`.
-.. toctree::
- :hidden:
- :maxdepth: 1
+```{toctree}
+:hidden: true
+:maxdepth: 1
- design/Overview
+design/Overview
+```
-OpenACC Support
-===============
+# OpenACC Support
-:doc:`OpenACC support <openacc/Overview>` is under development for
-both Flang and Clang. For this purpose, LLVM's OpenMP runtimes are
-being extended to serve as OpenACC runtimes. In some cases, Clang
-supports :doc:`OpenMP extensions <openacc/OpenMPExtensions>` to make
+{doc}`OpenACC support <openacc/Overview>` is under development for
+both Flang and Clang. For this purpose, LLVM's OpenMP runtimes are
+being extended to serve as OpenACC runtimes. In some cases, Clang
+supports {doc}`OpenMP extensions <openacc/OpenMPExtensions>` to make
the additional functionality also available in OpenMP applications.
-.. toctree::
- :hidden:
- :maxdepth: 1
+```{toctree}
+:hidden: true
+:maxdepth: 1
- openacc/Overview
+openacc/Overview
+```
-LLVM/OpenMP Optimizations
-=========================
+# LLVM/OpenMP Optimizations
-LLVM, since `version 11 <https://releases.llvm.org/download.html#11.0.0>`_ (12 Oct
-2020), has an :doc:`OpenMP-Aware optimization pass <optimizations/OpenMPOpt>`
-as well as the ability to :doc:`perform "scalar optimizations" across OpenMP region
+LLVM, since [version 11](https://releases.llvm.org/download.html#11.0.0) (12 Oct
+2020), has an {doc}`OpenMP-Aware optimization pass <optimizations/OpenMPOpt>`
+as well as the ability to {doc}`perform "scalar optimizations" across OpenMP region
boundaries <optimizations/OpenMPUnawareOptimizations>`.
-In-depth discussion of the topic can be found :doc:`here <optimizations/Overview>`.
+In-depth discussion of the topic can be found {doc}`here <optimizations/Overview>`.
-.. toctree::
- :hidden:
- :maxdepth: 1
+```{toctree}
+:hidden: true
+:maxdepth: 1
- optimizations/Overview
+optimizations/Overview
+```
-LLVM/OpenMP Optimization Remarks
-================================
+# LLVM/OpenMP Optimization Remarks
-LLVM has an elaborate ecosystem around `analysis and optimization remarks
-<https://llvm.org/docs/Remarks.html>`_ issues during
-compilation. The remarks can be enabled from the clang frontend `[1]`_ `[2]`_
-in various formats `[3]`_ `[4]`_ to be used by tools, i.a., `opt-viewer` or
+LLVM has an elaborate ecosystem around [analysis and optimization remarks](https://llvm.org/docs/Remarks.html) issues during
+compilation. The remarks can be enabled from the clang frontend [[1]][[1]] [[2]][[2]]
+in various formats [[3]][[3]] [[4]][[4]] to be used by tools, i.a., `opt-viewer` or
`llvm-opt-report` (dated).
The OpenMP optimizations in LLVM have been developed with remark support as a
priority. For a list of OpenMP specific remarks and more information on them,
-please refer to :doc:`remarks/OptimizationRemarks`.
-
+please refer to {doc}`remarks/OptimizationRemarks`.
-.. _`[1]`: https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports
-.. _`[2]`: https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags
-.. _`[3]`: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file
-.. _`[4]`: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record
+- [[1]][[1]] <https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports>
+- [[2]][[2]] <https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags>
+- [[3]][[3]] <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file>
+- [[4]][[4]] <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record>
-+ `[1]`_ https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports
-+ `[2]`_ https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags
-+ `[3]`_ https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file
-+ `[4]`_ https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record
+```{toctree}
+:hidden: true
+:maxdepth: 1
+remarks/OptimizationRemarks
+```
-.. toctree::
- :hidden:
- :maxdepth: 1
+# OpenMP Command-Line Argument Reference
- remarks/OptimizationRemarks
-
-OpenMP Command-Line Argument Reference
-======================================
-In addition to the
-`Clang command-line argument reference <https://clang.llvm.org/docs/ClangCommandLineReference.html>`_
-we also recommend the OpenMP
-:doc:`command-line argument reference <CommandLineArgumentReference>`
-page that offers a detailed overview of options specific to OpenMP. It also
+In addition to the
+[Clang command-line argument reference](https://clang.llvm.org/docs/ClangCommandLineReference.html)
+we also recommend the OpenMP
+{doc}`command-line argument reference <CommandLineArgumentReference>`
+page that offers a detailed overview of options specific to OpenMP. It also
contains a list of OpenMP offloading related command-line arguments.
+```{toctree}
+:hidden: true
+:maxdepth: 1
-.. toctree::
- :hidden:
- :maxdepth: 1
-
- CommandLineArgumentReference
+CommandLineArgumentReference
+```
-Support, Getting Involved, and Frequently Asked Questions (FAQ)
-===============================================================
+# Support, Getting Involved, and Frequently Asked Questions (FAQ)
Dealing with OpenMP can be complicated. For help with the setup of an OpenMP
(offload) capable compiler toolchain, its usage, and common problems, consult
-the :doc:`Support and FAQ <SupportAndFAQ>` page.
+the {doc}`Support and FAQ <SupportAndFAQ>` page.
-We also encourage everyone interested in OpenMP in LLVM to :doc:`get involved
+We also encourage everyone interested in OpenMP in LLVM to {doc}`get involved
<SupportAndFAQ>`.
+```{toctree}
+:hidden: true
+:maxdepth: 1
-.. toctree::
- :hidden:
- :maxdepth: 1
+SupportAndFAQ
+```
- SupportAndFAQ
+# Release Notes
-Release Notes
-=============
+The current (in-progress) release notes can be found {doc}`here <ReleaseNotes>` while
+release notes for releases, starting with LLVM 12, will be available on [the
+Download Page](https://releases.llvm.org/download.html).
-The current (in-progress) release notes can be found :doc:`here <ReleaseNotes>` while
-release notes for releases, starting with LLVM 12, will be available on `the
-Download Page <https://releases.llvm.org/download.html>`_.
+```{toctree}
+:hidden: true
+:maxdepth: 1
+In-Progress ReleaseNotes <ReleaseNotes>
+```
-.. toctree::
- :hidden:
- :maxdepth: 1
+[[1]]: https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports
+[[2]]: https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags
+[[3]]: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file
+[[4]]: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record
- In-Progress ReleaseNotes <ReleaseNotes>
diff --git a/openmp/docs/openacc/OpenMPExtensions.md b/openmp/docs/openacc/OpenMPExtensions.md
index 82cf77548b9a00..08516eb05ffff9 100644
--- a/openmp/docs/openacc/OpenMPExtensions.md
+++ b/openmp/docs/openacc/OpenMPExtensions.md
@@ -1,172 +1,163 @@
-OpenMP Extensions for OpenACC
-=============================
+# OpenMP Extensions for OpenACC
-OpenACC provides some functionality that OpenMP does not. In some
+OpenACC provides some functionality that OpenMP does not. In some
cases, Clang supports OpenMP extensions to provide similar
functionality, taking advantage of the runtime implementation already
-required for OpenACC. This section documents those extensions.
+required for OpenACC. This section documents those extensions.
-By default, Clang recognizes these extensions. The command-line
-option ``-fno-openmp-extensions`` can be specified to disable all
+By default, Clang recognizes these extensions. The command-line
+option `-fno-openmp-extensions` can be specified to disable all
OpenMP extensions, including those described in this section.
-.. _ompx-motivation:
+(ompx-motivation)=
-Motivation
-----------
+## Motivation
There are multiple benefits to exposing OpenACC functionality as LLVM
OpenMP extensions:
-* OpenMP applications can take advantage of the additional
+- OpenMP applications can take advantage of the additional
functionality.
-* As LLVM's implementation of these extensions matures, it can serve
+- As LLVM's implementation of these extensions matures, it can serve
as a basis for including these extensions in the OpenMP standard.
-* Source-to-source translation from certain OpenACC features to OpenMP
+- Source-to-source translation from certain OpenACC features to OpenMP
is otherwise impossible.
-* Runtime tests can be written in terms of OpenMP instead of OpenACC
+- Runtime tests can be written in terms of OpenMP instead of OpenACC
or low-level runtime calls.
-* More generally, there is a clean separation of concerns between
- OpenACC and OpenMP development in LLVM. That is, LLVM's OpenMP
+- More generally, there is a clean separation of concerns between
+ OpenACC and OpenMP development in LLVM. That is, LLVM's OpenMP
developers can discuss, modify, and debug LLVM's extended OpenMP
implementation and test suite without directly considering OpenACC's
language and execution model, which are handled by LLVM's OpenACC
developers.
-.. _ompx-hold:
+(ompx-hold)=
-``ompx_hold`` Map Type Modifier
--------------------------------
+## `ompx_hold` Map Type Modifier
-.. _ompx-holdExample:
+(ompx-holdexample)=
-Example
-^^^^^^^
+### Example
-.. code-block:: c++
+```c++
+#pragma omp target data map(ompx_hold, tofrom: x) // holds onto mapping of x throughout region
+{
+ foo(); // might have map(delete: x)
+ #pragma omp target map(present, alloc: x) // x is guaranteed to be present
+ printf("%d\n", x);
+}
+```
- #pragma omp target data map(ompx_hold, tofrom: x) // holds onto mapping of x throughout region
- {
- foo(); // might have map(delete: x)
- #pragma omp target map(present, alloc: x) // x is guaranteed to be present
- printf("%d\n", x);
- }
+The `ompx_hold` map type modifier above specifies that the `target
+data` directive holds onto the mapping for `x` throughout the
+associated region regardless of any `target exit data` directives
+executed during the call to `foo`. Thus, the presence assertion for
+`x` at the enclosed `target` construct cannot fail.
-The ``ompx_hold`` map type modifier above specifies that the ``target
-data`` directive holds onto the mapping for ``x`` throughout the
-associated region regardless of any ``target exit data`` directives
-executed during the call to ``foo``. Thus, the presence assertion for
-``x`` at the enclosed ``target`` construct cannot fail.
+(ompx-holdbehavior)=
-.. _ompx-holdBehavior:
+### Behavior
-Behavior
-^^^^^^^^
-
-* Stated more generally, the ``ompx_hold`` map type modifier specifies
+- Stated more generally, the `ompx_hold` map type modifier specifies
that the associated data is not unmapped until the end of the
- construct. As usual, the standard OpenMP reference count for the
+ construct. As usual, the standard OpenMP reference count for the
data must also reach zero before the data is unmapped.
-* If ``ompx_hold`` is specified for the same data on lexically or
+- If `ompx_hold` is specified for the same data on lexically or
dynamically enclosed constructs, there is no additional effect as
the data mapping is already held throughout their regions.
-* The ``ompx_hold`` map type modifier is permitted to appear only on
- ``target`` constructs (and associated combined constructs) and
- ``target data`` constructs. It is not permitted to appear on
- ``target enter data`` or ``target exit data`` directives because
+- The `ompx_hold` map type modifier is permitted to appear only on
+ `target` constructs (and associated combined constructs) and
+ `target data` constructs. It is not permitted to appear on
+ `target enter data` or `target exit data` directives because
there is no associated statement, so it is not meaningful to hold
onto a mapping until the end of the directive.
-* The runtime reports an error if ``omp_target_disassociate_ptr`` is
- called for a mapping for which the ``ompx_hold`` map type modifier
+- The runtime reports an error if `omp_target_disassociate_ptr` is
+ called for a mapping for which the `ompx_hold` map type modifier
is in effect.
-* Like the ``present`` map type modifier, the ``ompx_hold`` map type
+- Like the `present` map type modifier, the `ompx_hold` map type
modifier applies to an entire struct if it's specified for any
- member of that struct even if other ``map`` clauses on the same
- directive specify other members without the ``ompx_hold`` map type
+ member of that struct even if other `map` clauses on the same
+ directive specify other members without the `ompx_hold` map type
modifier.
-* ``ompx_hold`` support is not yet provided for ``defaultmap``.
+- `ompx_hold` support is not yet provided for `defaultmap`.
-Implementation
-^^^^^^^^^^^^^^
+### Implementation
-* LLVM uses the term *dynamic reference count* for the standard OpenMP
+- LLVM uses the term *dynamic reference count* for the standard OpenMP
reference count for host/device data mappings.
-* The ``ompx_hold`` map type modifier selects an alternate reference
+- The `ompx_hold` map type modifier selects an alternate reference
count, called the *hold reference count*.
-* A mapping is removed only once both its reference counts reach zero.
-* Because ``ompx_hold`` can appear only constructs, increments and
+- A mapping is removed only once both its reference counts reach zero.
+- Because `ompx_hold` can appear only constructs, increments and
decrements of the hold reference count are guaranteed to be
balanced, so it is impossible to decrement it below zero.
-* The dynamic reference count is used wherever ``ompx_hold`` is not
- specified (and possibly cannot be specified). Decrementing the
+- The dynamic reference count is used wherever `ompx_hold` is not
+ specified (and possibly cannot be specified). Decrementing the
dynamic reference count has no effect if it is already zero.
-* The runtime determines that the ``ompx_hold`` map type modifier is
- *in effect* (see :ref:`Behavior <ompx-holdBehavior>` above) when the
+- The runtime determines that the `ompx_hold` map type modifier is
+ *in effect* (see {ref}`Behavior <ompx-holdBehavior>` above) when the
hold reference count is greater than zero.
-Relationship with OpenACC
-^^^^^^^^^^^^^^^^^^^^^^^^^
+### Relationship with OpenACC
OpenACC specifies two reference counts for tracking host/device data
-mappings. Which reference count is used to implement an OpenACC
+mappings. Which reference count is used to implement an OpenACC
directive is determined by the nature of that directive, either
dynamic or structured:
-* The *dynamic reference count* is always used for ``enter data`` and
- ``exit data`` directives and corresponding OpenACC routines.
-* The *structured reference count* is always used for ``data`` and
- compute constructs, which are similar to OpenMP's ``target data``
- and ``target`` constructs.
+- The *dynamic reference count* is always used for `enter data` and
+ `exit data` directives and corresponding OpenACC routines.
+- The *structured reference count* is always used for `data` and
+ compute constructs, which are similar to OpenMP's `target data`
+ and `target` constructs.
Contrast with OpenMP, where the dynamic reference count is always used
unless the application developer specifies an alternate behavior via
-our map type modifier extension. We chose the name *hold* for that
-map type modifier because, as demonstrated in the above :ref:`example
+our map type modifier extension. We chose the name *hold* for that
+map type modifier because, as demonstrated in the above {ref}`example
<ompx-holdExample>`, *hold* concisely identifies the desired behavior
from the application developer's perspective without referencing the
implementation of that behavior.
The hold reference count is otherwise modeled after OpenACC's
-structured reference count. For example, calling ``acc_unmap_data``,
-which is similar to ``omp_target_disassociate_ptr``, is an error when
+structured reference count. For example, calling `acc_unmap_data`,
+which is similar to `omp_target_disassociate_ptr`, is an error when
the structured reference count is not zero.
While Flang and Clang obviously must implement the syntax and
semantics for selecting OpenACC reference counts differently than for
selecting OpenMP reference counts, the implementation is the same at
-the runtime level. That is, OpenACC's dynamic reference count is
+the runtime level. That is, OpenACC's dynamic reference count is
OpenMP's dynamic reference count, and OpenACC's structured reference
count is our OpenMP hold reference count extension.
-.. _atomicWithinTeams:
+(atomicwithinteams)=
-``atomic`` Strictly Nested Within ``teams``
--------------------------------------------
+## `atomic` Strictly Nested Within `teams`
-Example
-^^^^^^^
+### Example
OpenMP 5.2, sec. 10.2 "teams Construct", p. 232, L9-12 restricts what
-regions can be strictly nested within a ``teams`` region. As an
+regions can be strictly nested within a `teams` region. As an
extension, Clang relaxes that restriction in the case of the
-``atomic`` construct so that, for example, the following case is
+`atomic` construct so that, for example, the following case is
permitted:
-.. code-block:: c++
-
- #pragma omp target teams map(tofrom:x)
- #pragma omp atomic update
- x++;
+```c++
+#pragma omp target teams map(tofrom:x)
+#pragma omp atomic update
+x++;
+```
-Relationship with OpenACC
-^^^^^^^^^^^^^^^^^^^^^^^^^
+### Relationship with OpenACC
This extension is important when translating OpenACC to OpenMP because
OpenACC does not have the same restriction for its corresponding
-constructs. For example, the following is conforming OpenACC:
+constructs. For example, the following is conforming OpenACC:
-.. code-block:: c++
+```c++
+#pragma acc parallel copy(x)
+#pragma acc atomic update
+x++;
+```
- #pragma acc parallel copy(x)
- #pragma acc atomic update
- x++;
diff --git a/openmp/docs/openacc/Overview.md b/openmp/docs/openacc/Overview.md
index 6dbf36201951e1..2d466428394958 100644
--- a/openmp/docs/openacc/Overview.md
+++ b/openmp/docs/openacc/Overview.md
@@ -1,13 +1,14 @@
-OpenACC Support
-===============
+# OpenACC Support
-OpenACC support is under development for both Flang and Clang. For
+OpenACC support is under development for both Flang and Clang. For
this purpose, LLVM's OpenMP runtimes are being extended to serve as
OpenACC runtimes.
-.. toctree::
- :glob:
- :hidden:
- :maxdepth: 1
+```{toctree}
+:glob: true
+:hidden: true
+:maxdepth: 1
+
+OpenMPExtensions
+```
- OpenMPExtensions
diff --git a/openmp/docs/optimizations/OpenMPOpt.md b/openmp/docs/optimizations/OpenMPOpt.md
index 8de9e21ad1ff27..8e9924da2138b0 100644
--- a/openmp/docs/optimizations/OpenMPOpt.md
+++ b/openmp/docs/optimizations/OpenMPOpt.md
@@ -1,33 +1,30 @@
-==========================
-OpenMP-Aware Optimizations
-==========================
+# OpenMP-Aware Optimizations
-LLVM, since `version 11 <https://releases.llvm.org/download.html#11.0.0>`_ (12
-Oct 2020), supports an :ref:`OpenMP-Aware optimization pass <OpenMPOpt>`. This
+LLVM, since [version 11](https://releases.llvm.org/download.html#11.0.0) (12
+Oct 2020), supports an {ref}`OpenMP-Aware optimization pass <OpenMPOpt>`. This
optimization pass will attempt to optimize the module with OpenMP-specific
domain-knowledge. This pass is enabled by default at high optimization levels
(O2 / O3) if compiling with OpenMP support enabled.
-.. _OpenMPOpt:
+(openmpopt)=
-OpenMPOpt
-=========
+## OpenMPOpt
-.. contents::
- :local:
- :depth: 1
+```{contents}
+:depth: 1
+:local: true
+```
OpenMPOpt contains several OpenMP-Aware optimizations. This pass is run early on
the entire Module, and later on the entire call graph. Most optimizations done
by OpenMPOpt support remarks. Optimization remarks can be enabled by compiling
with the following flags.
-.. code-block:: console
+```console
+$ clang -Rpass=openmp-opt -Rpass-missed=openmp-opt -Rpass-analysis=openmp-opt
+```
- $ clang -Rpass=openmp-opt -Rpass-missed=openmp-opt -Rpass-analysis=openmp-opt
-
-OpenMP Runtime Call Deduplication
----------------------------------
+### OpenMP Runtime Call Deduplication
The OpenMP runtime library contains several functions used to implement features
of the OpenMP standard. Several of the runtime calls are constant within a
@@ -36,8 +33,7 @@ single reference, but in this case the compiler will only see an opaque call
into the runtime library. To get around this, OpenMPOpt maintains a list of
OpenMP runtime functions that are constant and will manually deduplicate them.
-Globalization
--------------
+### Globalization
The OpenMP standard requires that data can be shared between different threads.
This requirement poses a unique challenge when offloading to GPU accelerators.
@@ -50,19 +46,19 @@ generating code for this offloading region, it will see that the variable `x`
escapes and is potentially shared. This will require globalizing the variable,
which means it cannot reside in the registers on the device.
-.. code-block:: c++
-
- void use(void *) { }
+```c++
+void use(void *) { }
- void foo() {
- int x;
- use(&x);
- }
+void foo() {
+ int x;
+ use(&x);
+}
- int main() {
- #pragma omp target parallel
- foo();
- }
+int main() {
+#pragma omp target parallel
+ foo();
+}
+```
In many cases, this transformation is not actually necessary but still carries a
significant performance penalty. Because of this, OpenMPOpt can perform and
@@ -73,25 +69,24 @@ register memory.
Another case is memory that is intentionally shared between the threads, but is
shared from one thread to all the others. Such variables can be moved to shared
-memory when compiled without needing to go through the runtime library. This
+memory when compiled without needing to go through the runtime library. This
allows for users to confidently declare shared memory on the device without
needing to use custom OpenMP allocators or rely on the runtime.
+```c++
+static void share(void *);
-.. code-block:: c++
-
- static void share(void *);
+static void foo() {
+ int x[64];
+#pragma omp parallel
+ share(x);
+}
- static void foo() {
- int x[64];
- #pragma omp parallel
- share(x);
- }
-
- int main() {
- #pragma omp target
- foo();
- }
+int main() {
+ #pragma omp target
+ foo();
+}
+```
These optimizations can have very large performance implications. Both of these
optimizations rely heavily on inter-procedural analysis. Because of this,
@@ -100,9 +95,9 @@ and functions should not be externally visible unless needed. OpenMPOpt will
inform the user if any globalization calls remain if remarks are enabled. This
should be treated as a defect in the program.
-Resources
-=========
+## Resources
+
+- 2021 OpenMP Webinar: "A Compiler's View of OpenMP" <https://youtu.be/eIMpgez61r4>
+- 2020 LLVM Developers’ Meeting: "(OpenMP) Parallelism-Aware Optimizations" <https://youtu.be/gtxWkeLCxmU>
+- 2019 EuroLLVM Developers’ Meeting: "Compiler Optimizations for (OpenMP) Target Offloading to GPUs" <https://youtu.be/3AbS82C3X30>
-- 2021 OpenMP Webinar: "A Compiler's View of OpenMP" https://youtu.be/eIMpgez61r4
-- 2020 LLVM Developers’ Meeting: "(OpenMP) Parallelism-Aware Optimizations" https://youtu.be/gtxWkeLCxmU
-- 2019 EuroLLVM Developers’ Meeting: "Compiler Optimizations for (OpenMP) Target Offloading to GPUs" https://youtu.be/3AbS82C3X30
diff --git a/openmp/docs/optimizations/OpenMPUnawareOptimizations.md b/openmp/docs/optimizations/OpenMPUnawareOptimizations.md
index 72ccc0355273ee..60447cceacbb7e 100644
--- a/openmp/docs/optimizations/OpenMPUnawareOptimizations.md
+++ b/openmp/docs/optimizations/OpenMPUnawareOptimizations.md
@@ -1,9 +1,7 @@
-OpenMP-Unaware Optimizations
-============================
+# OpenMP-Unaware Optimizations
+## Resources
-Resources
----------
+- 2018 LLVM Developers’ Meeting: "Optimizing Indirections, using abstractions without remorse" <https://youtu.be/zfiHaPaoQPc>
+- 2019 LLVM Developers’ Meeting: "The Attributor: A Versatile Inter-procedural Fixpoint Iteration Framework" <https://youtu.be/CzWkc_JcfS0>
-- 2018 LLVM Developers’ Meeting: "Optimizing Indirections, using abstractions without remorse" https://youtu.be/zfiHaPaoQPc
-- 2019 LLVM Developers’ Meeting: "The Attributor: A Versatile Inter-procedural Fixpoint Iteration Framework" https://youtu.be/CzWkc_JcfS0
diff --git a/openmp/docs/optimizations/Overview.md b/openmp/docs/optimizations/Overview.md
index 966757d206fdeb..4559b448704abb 100644
--- a/openmp/docs/optimizations/Overview.md
+++ b/openmp/docs/optimizations/Overview.md
@@ -1,19 +1,19 @@
-.. _llvm_openmp_optimizations:
+(llvm-openmp-optimizations)=
-OpenMP Optimizations in LLVM
-============================
+# OpenMP Optimizations in LLVM
-LLVM, since `version 11 <https://releases.llvm.org/download.html#11.0.0>`_ (12 Oct
-2020), has an :doc:`OpenMP-Aware optimization pass <OpenMPOpt>`
-as well as the ability to :doc:`perform "scalar optimizations" across OpenMP region
+LLVM, since [version 11](https://releases.llvm.org/download.html#11.0.0) (12 Oct
+2020), has an {doc}`OpenMP-Aware optimization pass <OpenMPOpt>`
+as well as the ability to {doc}`perform "scalar optimizations" across OpenMP region
boundaries <OpenMPUnawareOptimizations>`.
-.. toctree::
- :glob:
- :hidden:
- :titlesonly:
- :maxdepth: 1
+```{toctree}
+:glob: true
+:hidden: true
+:maxdepth: 1
+:titlesonly: true
- OpenMPOpt
- OpenMPUnawareOptimizations
+OpenMPOpt
+OpenMPUnawareOptimizations
+```
diff --git a/openmp/docs/remarks/OMP100.md b/openmp/docs/remarks/OMP100.md
index 30c9fec2f67b58..ce475d7487be80 100644
--- a/openmp/docs/remarks/OMP100.md
+++ b/openmp/docs/remarks/OMP100.md
@@ -1,8 +1,8 @@
-.. _omp100:
-.. _omp_no_external_caller_in_target_region:
+(omp100)=
-Potentially unknown OpenMP target region caller `[OMP100]`
-==========================================================
+(omp-no-external-caller-in-target-region)=
+
+# Potentially unknown OpenMP target region caller `[OMP100]`
A function remark that indicates the function, when compiled for a GPU, is
potentially called from outside the translation unit. Note that a remark is
@@ -24,3 +24,4 @@ impossible, e.g., because there are outside callers on the host, one can split
the function into an external visible interface which is not compiled for
the target and an internal implementation which is compiled for the target
and should be called from within the target region.
+
diff --git a/openmp/docs/remarks/OMP101.md b/openmp/docs/remarks/OMP101.md
index b7245c68c5ac3d..79b55799506d0a 100644
--- a/openmp/docs/remarks/OMP101.md
+++ b/openmp/docs/remarks/OMP101.md
@@ -1,6 +1,6 @@
-.. _omp101:
+(omp101)=
-Parallel region is used in unknown / unexpected ways. Will not attempt to rewrite the state machine. [OMP101]
-=============================================================================================================
+# Parallel region is used in unknown / unexpected ways. Will not attempt to rewrite the state machine. [OMP101]
An analysis remark that indicates that a parallel region has unknown calls.
+
diff --git a/openmp/docs/remarks/OMP102.md b/openmp/docs/remarks/OMP102.md
index c8d37bc283adf2..d4e659216d66c9 100644
--- a/openmp/docs/remarks/OMP102.md
+++ b/openmp/docs/remarks/OMP102.md
@@ -1,8 +1,8 @@
-.. _omp102:
+(omp102)=
-Parallel region is not called from a unique kernel. Will not attempt to rewrite the state machine. [OMP102]
-===========================================================================================================
+# Parallel region is not called from a unique kernel. Will not attempt to rewrite the state machine. [OMP102]
This analysis remark indicates that a given parallel region is called by
multiple kernels. This prevents the compiler from optimizing it to a single
kernel and rewrite the state machine.
+
diff --git a/openmp/docs/remarks/OMP110.md b/openmp/docs/remarks/OMP110.md
index ae01b4d9a41ab5..9fe27c3beabddd 100644
--- a/openmp/docs/remarks/OMP110.md
+++ b/openmp/docs/remarks/OMP110.md
@@ -1,7 +1,6 @@
-.. _omp110:
+(omp110)=
-Moving globalized variable to the stack. [OMP110]
-=================================================
+# Moving globalized variable to the stack. [OMP110]
This optimization remark indicates that a globalized variable was moved back to
thread-local stack memory on the device. This occurs when the optimization pass
@@ -20,64 +19,63 @@ that can be shared, such as shared or global memory. This optimization moves the
data back from shared or global memory to thread-local stack memory if the data
is not actually shared between the threads.
-Examples
---------
+## Examples
A trivial example of globalization occurring can be seen with this example. The
-compiler sees that a pointer to the thread-local variable ``x`` escapes the
+compiler sees that a pointer to the thread-local variable `x` escapes the
current scope and must globalize it even though it is not actually necessary.
Fortunately, this optimization can undo this by looking at its usage.
-.. code-block:: c++
+```c++
+void use(int *x) { }
- void use(int *x) { }
+void foo() {
+ int x;
+ use(&x);
+}
- void foo() {
- int x;
- use(&x);
- }
+int main() {
+#pragma omp target parallel
+ foo();
+}
+```
- int main() {
- #pragma omp target parallel
- foo();
- }
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 omp110.cpp -O1 -Rpass=openmp-opt
- omp110.cpp:6:7: remark: Moving globalized variable to the stack. [OMP110]
- int x;
- ^
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 omp110.cpp -O1 -Rpass=openmp-opt
+omp110.cpp:6:7: remark: Moving globalized variable to the stack. [OMP110]
+ int x;
+ ^
+```
A less trivial example can be seen using C++'s complex numbers. In this case the
overloaded arithmetic operators cause pointers to the complex numbers to escape
the current scope, but they can again be removed once the usage is visible.
-.. code-block:: c++
-
- #include <complex>
+```c++
+#include <complex>
- using complex = std::complex<double>;
+using complex = std::complex<double>;
- void zaxpy(complex *X, complex *Y, const complex D, int N) {
- #pragma omp target teams distribute parallel for firstprivate(D)
- for (int i = 0; i < N; ++i)
- Y[i] = D * X[i] + Y[i];
- }
+void zaxpy(complex *X, complex *Y, const complex D, int N) {
+#pragma omp target teams distribute parallel for firstprivate(D)
+ for (int i = 0; i < N; ++i)
+ Y[i] = D * X[i] + Y[i];
+}
+```
-.. code-block:: console
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 omp110.cpp -O1 -Rpass=openmp-opt
+In file included from omp110.cpp:1:
+In file included from /usr/bin/clang/lib/clang/13.0.0/include/openmp_wrappers/complex:27:
+/usr/include/c++/8/complex:328:20: remark: Moving globalized variable to the stack. [OMP110]
+ complex<_Tp> __r = __x;
+ ^
+/usr/include/c++/8/complex:388:20: remark: Moving globalized variable to the stack. [OMP110]
+ complex<_Tp> __r = __x;
+ ^
+```
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 omp110.cpp -O1 -Rpass=openmp-opt
- In file included from omp110.cpp:1:
- In file included from /usr/bin/clang/lib/clang/13.0.0/include/openmp_wrappers/complex:27:
- /usr/include/c++/8/complex:328:20: remark: Moving globalized variable to the stack. [OMP110]
- complex<_Tp> __r = __x;
- ^
- /usr/include/c++/8/complex:388:20: remark: Moving globalized variable to the stack. [OMP110]
- complex<_Tp> __r = __x;
- ^
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading optimization remark.
+
diff --git a/openmp/docs/remarks/OMP111.md b/openmp/docs/remarks/OMP111.md
index d7b87bd9ea8ac3..c58b82de48e55c 100644
--- a/openmp/docs/remarks/OMP111.md
+++ b/openmp/docs/remarks/OMP111.md
@@ -1,7 +1,6 @@
-.. _omp111:
+(omp111)=
-Replaced globalized variable with X bytes of shared memory. [OMP111]
-====================================================================
+# Replaced globalized variable with X bytes of shared memory. [OMP111]
This optimization occurs when a globalized variable's data is shared between
multiple threads, but requires a constant amount of memory that can be
@@ -19,56 +18,54 @@ that the parent functions have internal linkage. Otherwise, an external caller
could invalidate this analysis but having multiple threads call that function.
The optimization pass will make internal copies of each function to use for this
reason, but it is still recommended to mark them as internal using keywords like
-``static`` whenever possible.
+`static` whenever possible.
-Example
--------
+## Example
This optimization should apply to any variable declared in an OpenMP target
region that is then shared with every thread in a parallel region. This allows
the user to declare shared memory without using custom allocators. A simple
stencil calculation shows how this can be used.
-.. code-block:: c++
-
- void stencil(int M, int N, double *X, double *Y) {
- #pragma omp target teams distribute collapse(2) \
- map(to : X [0:M * N]) map(tofrom : Y [0:M * N])
- for (int i0 = 0; i0 < M; i0 += MC) {
- for (int j0 = 0; j0 < N; j0 += NC) {
- double sX[MC][NC];
+```c++
+void stencil(int M, int N, double *X, double *Y) {
+#pragma omp target teams distribute collapse(2) \
+ map(to : X [0:M * N]) map(tofrom : Y [0:M * N])
+ for (int i0 = 0; i0 < M; i0 += MC) {
+ for (int j0 = 0; j0 < N; j0 += NC) {
+ double sX[MC][NC];
- #pragma omp parallel for collapse(2) shared(sX) default(firstprivate)
- for (int i1 = 0; i1 < MC; ++i1)
- for (int j1 = 0; j1 < NC; ++j1)
- sX[i1][j1] = X[(i0 + i1) * N + (j0 + j1)];
+#pragma omp parallel for collapse(2) shared(sX) default(firstprivate)
+ for (int i1 = 0; i1 < MC; ++i1)
+ for (int j1 = 0; j1 < NC; ++j1)
+ sX[i1][j1] = X[(i0 + i1) * N + (j0 + j1)];
- #pragma omp parallel for collapse(2) shared(sX) default(firstprivate)
- for (int i1 = 1; i1 < MC - 1; ++i1)
- for (int j1 = 1; j1 < NC - 1; ++j1)
- Y[(i0 + i1) * N + j0 * j1] = (sX[i1 + 1][j1] + sX[i1 - 1][j1] +
- sX[i1][j1 + 1] + sX[i1][j1 - 1] +
- -4.0 * sX[i1][j1]) / (dX * dX);
- }
+#pragma omp parallel for collapse(2) shared(sX) default(firstprivate)
+ for (int i1 = 1; i1 < MC - 1; ++i1)
+ for (int j1 = 1; j1 < NC - 1; ++j1)
+ Y[(i0 + i1) * N + j0 * j1] = (sX[i1 + 1][j1] + sX[i1 - 1][j1] +
+ sX[i1][j1 + 1] + sX[i1][j1 - 1] +
+ -4.0 * sX[i1][j1]) / (dX * dX);
}
}
+}
+```
-.. code-block:: console
-
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass=openmp-opt -fopenmp-version=51 omp111.cpp
- omp111.cpp:10:14: remark: Replaced globalized variable with 8192 bytes of shared memory. [OMP111]
- double sX[MC][NC];
- ^
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass=openmp-opt -fopenmp-version=51 omp111.cpp
+omp111.cpp:10:14: remark: Replaced globalized variable with 8192 bytes of shared memory. [OMP111]
+ double sX[MC][NC];
+ ^
+```
The default mapping for variables captured in an OpenMP parallel region is
-``shared``. This means taking a pointer to the object which will ultimately
+`shared`. This means taking a pointer to the object which will ultimately
result in globalization that will be mapped to shared memory when it could have
been placed in registers. To avoid this, make sure each variable that can be
-copied into the region is marked ``firstprivate`` either explicitly or using the
-OpenMP 5.1 feature ``default(firstprivate)``.
+copied into the region is marked `firstprivate` either explicitly or using the
+OpenMP 5.1 feature `default(firstprivate)`.
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading optimization remark.
+
diff --git a/openmp/docs/remarks/OMP112.md b/openmp/docs/remarks/OMP112.md
index a992b4e41969f6..84fbe2488dcd05 100644
--- a/openmp/docs/remarks/OMP112.md
+++ b/openmp/docs/remarks/OMP112.md
@@ -1,11 +1,10 @@
-.. _omp112:
+(omp112)=
-Found thread data sharing on the GPU. Expect degraded performance due to data globalization. [OMP112]
-=====================================================================================================
+# Found thread data sharing on the GPU. Expect degraded performance due to data globalization. [OMP112]
This missed remark indicates that a globalized value was found on the target
-device that was not either replaced with stack memory by :ref:`OMP110 <omp110>`
-or shared memory by :ref:`OMP111 <omp111>`. Globalization that has not been
+device that was not either replaced with stack memory by {ref}`OMP110 <omp110>`
+or shared memory by {ref}`OMP111 <omp111>`. Globalization that has not been
removed will need to be handled by the runtime and will significantly impact
performance.
@@ -18,72 +17,71 @@ shared between the threads. In the majority of cases, globalized variables can
either be returns to a thread-local stack, or pushed to shared memory. However,
in a few cases it is necessary and will cause a performance penalty.
-Examples
---------
+## Examples
This example shows legitimate data sharing on the device. It is a convoluted
example, but is completely complaint with the OpenMP standard. If globalization
was not added this would result in different results on different target
devices.
-.. code-block:: c++
+```c++
+#include <omp.h>
+#include <cstdio>
- #include <omp.h>
- #include <cstdio>
+#pragma omp declare target
+static int *p;
+#pragma omp end declare target
- #pragma omp declare target
- static int *p;
- #pragma omp end declare target
-
- void foo() {
- int x = omp_get_thread_num();
- if (omp_get_thread_num() == 1)
- p = &x;
+void foo() {
+ int x = omp_get_thread_num();
+ if (omp_get_thread_num() == 1)
+ p = &x;
- #pragma omp barrier
+#pragma omp barrier
- printf ("Thread %d: %d\n", omp_get_thread_num(), *p);
- }
+ printf ("Thread %d: %d\n", omp_get_thread_num(), *p);
+}
- int main() {
- #pragma omp target parallel
- foo();
- }
+int main() {
+#pragma omp target parallel
+ foo();
+}
+```
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-missed=openmp-opt omp112.cpp
- omp112.cpp:9:7: remark: Found thread data sharing on the GPU. Expect degraded performance
- due to data globalization. [OMP112] [-Rpass-missed=openmp-opt]
- int x = omp_get_thread_num();
- ^
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-missed=openmp-opt omp112.cpp
+omp112.cpp:9:7: remark: Found thread data sharing on the GPU. Expect degraded performance
+due to data globalization. [OMP112] [-Rpass-missed=openmp-opt]
+int x = omp_get_thread_num();
+ ^
+```
A less convoluted example globalization that cannot be removed occurs when
calling functions that aren't visible from the current translation unit.
-.. code-block:: c++
-
- extern void use(int *x);
+```c++
+extern void use(int *x);
- void foo() {
- int x;
- use(&x);
- }
-
- int main() {
- #pragma omp target parallel
- foo();
- }
+void foo() {
+ int x;
+ use(&x);
+}
-.. code-block:: console
+int main() {
+#pragma omp target parallel
+ foo();
+}
+```
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-missed=openmp-opt omp112.cpp
- omp112.cpp:4:7: remark: Found thread data sharing on the GPU. Expect degraded performance
- due to data globalization. [OMP112] [-Rpass-missed=openmp-opt]
- int x;
- ^
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-missed=openmp-opt omp112.cpp
+omp112.cpp:4:7: remark: Found thread data sharing on the GPU. Expect degraded performance
+due to data globalization. [OMP112] [-Rpass-missed=openmp-opt]
+int x;
+ ^
+```
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading missed remark.
+
diff --git a/openmp/docs/remarks/OMP113.md b/openmp/docs/remarks/OMP113.md
index ae826300057b55..9af7c3a3178623 100644
--- a/openmp/docs/remarks/OMP113.md
+++ b/openmp/docs/remarks/OMP113.md
@@ -1,7 +1,6 @@
-.. _omp113:
+(omp113)=
-Could not move globalized variable to the stack. Variable is potentially captured in call. Mark parameter as `__attribute__((noescape))` to override. [OMP113]
-==============================================================================================================================================================
+# Could not move globalized variable to the stack. Variable is potentially captured in call. Mark parameter as `__attribute__((noescape))` to override. [OMP113]
This missed remark indicates that a globalized value could not be moved to the
stack because it is potentially captured by a call to a function we cannot
@@ -16,66 +15,65 @@ shared if a copy of its pointer is never made. However, this remark indicates a
copy of the pointer is present or that sharing is possible because it is used
outside the current translation unit.
-Examples
---------
+## Examples
If a pointer to a thread-local variable is passed to a function not visible in
the current translation unit we need to assume a copy is made of it that can be
-shared between the threads. This prevents :ref:`OMP110 <omp110>` from
+shared between the threads. This prevents {ref}`OMP110 <omp110>` from
triggering, which will result in a performance penalty when executing on the
target device.
-.. code-block:: c++
-
- extern void use(int *x);
-
- void foo() {
- int x;
- use(&x);
- }
-
- int main() {
- #pragma omp target parallel
- foo();
- }
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-missed=openmp-opt omp113.cpp
- missed.cpp:4:7: remark: Could not move globalized variable to the stack. Variable is
- potentially captured in call. Mark parameter as `__attribute__((noescape))` to
- override. [OMP113]
- int x;
- ^
-
-As the remark suggests, this behaviour can be overridden using the ``noescape``
+```c++
+extern void use(int *x);
+
+void foo() {
+ int x;
+ use(&x);
+}
+
+int main() {
+#pragma omp target parallel
+ foo();
+}
+```
+
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-missed=openmp-opt omp113.cpp
+missed.cpp:4:7: remark: Could not move globalized variable to the stack. Variable is
+potentially captured in call. Mark parameter as `__attribute__((noescape))` to
+override. [OMP113]
+ int x;
+ ^
+```
+
+As the remark suggests, this behaviour can be overridden using the `noescape`
attribute. This tells the compiler that no reference to the object the pointer
points to that is derived from the parameter value will survive after the
function returns. The user is responsible for verifying that this assertion is
correct.
-.. code-block:: c++
-
- extern void use(__attribute__((noescape)) int *x);
+```c++
+extern void use(__attribute__((noescape)) int *x);
- void foo() {
- int x;
- use(&x);
- }
+void foo() {
+ int x;
+ use(&x);
+}
- int main() {
- #pragma omp target parallel
- foo();
- }
+int main() {
+#pragma omp target parallel
+ foo();
+}
+```
-.. code-block:: console
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp113.cpp
+missed.cpp:4:7: remark: Moving globalized variable to the stack. [OMP110]
+int x;
+ ^
+```
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp113.cpp
- missed.cpp:4:7: remark: Moving globalized variable to the stack. [OMP110]
- int x;
- ^
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading missed remark.
+
diff --git a/openmp/docs/remarks/OMP120.md b/openmp/docs/remarks/OMP120.md
index 1c3c45f8702e79..3b5a48e69f5b8f 100644
--- a/openmp/docs/remarks/OMP120.md
+++ b/openmp/docs/remarks/OMP120.md
@@ -1,7 +1,6 @@
-.. _omp120:
+(omp120)=
-Transformed generic-mode kernel to SPMD-mode [OMP120]
-=====================================================
+# Transformed generic-mode kernel to SPMD-mode [OMP120]
This optimization remark indicates that the execution strategy for the OpenMP
target offloading kernel was changed. Generic-mode kernels are executed by a
@@ -17,10 +16,9 @@ Generic-mode is often considerably slower than SPMD-mode because of the extra
overhead required to separately schedule worker threads and pass data between
them.This optimization allows users to use generic-mode semantics while
achieving the performance of SPMD-mode. This can be helpful when defining shared
-memory between the threads using :ref:`OMP111 <omp111>`.
+memory between the threads using {ref}`OMP111 <omp111>`.
-Examples
---------
+## Examples
Normally, any kernel that contains split OpenMP target and parallel regions will
be executed in generic-mode. Sometimes it is easier to use generic-mode
@@ -28,67 +26,66 @@ semantics to define shared memory, or more tightly control the distribution of
the threads. This shows a naive matrix-matrix multiplication that contains code
that will need to be guarded.
-.. code-block:: c++
+```c++
+void matmul(int M, int N, int K, double *A, double *B, double *C) {
+#pragma omp target teams distribute collapse(2) \
+ map(to:A[0: M*K]) map(to:B[0: K*N]) map(tofrom:C[0 : M*N])
+ for (int i = 0; i < M; i++) {
+ for (int j = 0; j < N; j++) {
+ double sum = 0.0;
- void matmul(int M, int N, int K, double *A, double *B, double *C) {
- #pragma omp target teams distribute collapse(2) \
- map(to:A[0: M*K]) map(to:B[0: K*N]) map(tofrom:C[0 : M*N])
- for (int i = 0; i < M; i++) {
- for (int j = 0; j < N; j++) {
- double sum = 0.0;
+#pragma omp parallel for reduction(+:sum) default(firstprivate)
+ for (int k = 0; k < K; k++)
+ sum += A[i*K + k] * B[k*N + j];
- #pragma omp parallel for reduction(+:sum) default(firstprivate)
- for (int k = 0; k < K; k++)
- sum += A[i*K + k] * B[k*N + j];
-
- C[i*N + j] = sum;
- }
+ C[i*N + j] = sum;
}
}
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -fopenmp-version=51 -O2 -Rpass=openmp-opt omp120.cpp
- omp120.cpp:6:14: remark: Replaced globalized variable with 8 bytes of shared memory. [OMP111]
- double sum = 0.0;
- ^
- omp120.cpp:2:1: remark: Transformed generic-mode kernel to SPMD-mode. [OMP120]
- #pragma omp target teams distribute collapse(2) \
- ^
-
-This requires guarding the store to the shared variable ``sum`` and the store to
-the matrix ``C``. This can be thought of as generating the code below.
-
-.. code-block:: c++
-
- void matmul(int M, int N, int K, double *A, double *B, double *C) {
- #pragma omp target teams distribute collapse(2) \
- map(to:A[0: M*K]) map(to:B[0: K*N]) map(tofrom:C[0 : M*N])
- for (int i = 0; i < M; i++) {
- for (int j = 0; j < N; j++) {
- double sum;
- #pragma omp parallel default(firstprivate) shared(sum)
- {
- #pragma omp barrier
- if (omp_get_thread_num() == 0)
- sum = 0.0;
- #pragma omp barrier
-
- #pragma omp for reduction(+:sum)
- for (int k = 0; k < K; k++)
- sum += A[i*K + k] * B[k*N + j];
-
- #pragma omp barrier
- if (omp_get_thread_num() == 0)
- C[i*N + j] = sum;
- #pragma omp barrier
- }
- }
+}
+```
+
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -fopenmp-version=51 -O2 -Rpass=openmp-opt omp120.cpp
+omp120.cpp:6:14: remark: Replaced globalized variable with 8 bytes of shared memory. [OMP111]
+ double sum = 0.0;
+ ^
+omp120.cpp:2:1: remark: Transformed generic-mode kernel to SPMD-mode. [OMP120]
+#pragma omp target teams distribute collapse(2) \
+^
+```
+
+This requires guarding the store to the shared variable `sum` and the store to
+the matrix `C`. This can be thought of as generating the code below.
+
+```c++
+void matmul(int M, int N, int K, double *A, double *B, double *C) {
+#pragma omp target teams distribute collapse(2) \
+ map(to:A[0: M*K]) map(to:B[0: K*N]) map(tofrom:C[0 : M*N])
+ for (int i = 0; i < M; i++) {
+ for (int j = 0; j < N; j++) {
+ double sum;
+#pragma omp parallel default(firstprivate) shared(sum)
+ {
+ #pragma omp barrier
+ if (omp_get_thread_num() == 0)
+ sum = 0.0;
+ #pragma omp barrier
+
+#pragma omp for reduction(+:sum)
+ for (int k = 0; k < K; k++)
+ sum += A[i*K + k] * B[k*N + j];
+
+ #pragma omp barrier
+ if (omp_get_thread_num() == 0)
+ C[i*N + j] = sum;
+ #pragma omp barrier
+ }
}
}
+}
+```
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading optimization remark.
+
diff --git a/openmp/docs/remarks/OMP121.md b/openmp/docs/remarks/OMP121.md
index f3ceeac7f3ab92..1df315abcd02e0 100644
--- a/openmp/docs/remarks/OMP121.md
+++ b/openmp/docs/remarks/OMP121.md
@@ -1,7 +1,6 @@
-.. _omp121:
+(omp121)=
-Value has potential side effects preventing SPMD-mode execution. Add `[[omp::assume(\"ompx_spmd_amenable\")]]` to the called function to override. [OMP121]
-===================================================================================================================================================================
+# Value has potential side effects preventing SPMD-mode execution. Add `[[omp::assume(\"ompx_spmd_amenable\")]]` to the called function to override. [OMP121]
This analysis remarks indicates that a potential side-effect that cannot be
guarded prevents the target region from executing in SPMD-mode. SPMD-mode
@@ -15,66 +14,64 @@ translation unit will prevent this transformation from occurring as well, but
can be overridden using an assumption stating that it contains no calls that
prevent SPMD execution.
-Examples
---------
+## Examples
Calls to functions outside the current translation unit may contain instructions
or operations that cannot be executed in SPMD-mode.
-.. code-block:: c++
+```c++
+extern int work();
- extern int work();
+void use(int x);
- void use(int x);
+void foo() {
+#pragma omp target teams
+ {
+ int x = work();
+#pragma omp parallel
+ use(x);
- void foo() {
- #pragma omp target teams
- {
- int x = work();
- #pragma omp parallel
- use(x);
-
- }
}
-
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-analysis=openmp-opt omp121.cpp
- omp121.cpp:8:13: remark: Value has potential side effects preventing SPMD-mode
- execution. Add `[[omp::assume("ompx_spmd_amenable")]]` to the called function
- to override. [OMP121]
- int x = work();
- ^
+}
+```
+
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-analysis=openmp-opt omp121.cpp
+omp121.cpp:8:13: remark: Value has potential side effects preventing SPMD-mode
+execution. Add `[[omp::assume("ompx_spmd_amenable")]]` to the called function
+to override. [OMP121]
+int x = work();
+ ^
+```
As the remark suggests, the problem is caused by the unknown call to the
-external function ``work``. This can be overridden by asserting that it does not
+external function `work`. This can be overridden by asserting that it does not
contain any code that prevents SPMD-mode execution.
-.. code-block:: c++
+```c++
+[[omp::assume("ompx_spmd_amenable")]] extern int work();
- [[omp::assume("ompx_spmd_amenable")]] extern int work();
+void use(int x);
- void use(int x);
+void foo() {
+#pragma omp target teams
+ {
+ int x = work();
+#pragma omp parallel
+ use(x);
- void foo() {
- #pragma omp target teams
- {
- int x = work();
- #pragma omp parallel
- use(x);
-
- }
}
+}
+```
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp121.cpp
- omp121.cpp:6:1: remark: Transformed generic-mode kernel to SPMD-mode. [OMP120]
- #pragma omp target teams
- ^
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp121.cpp
+omp121.cpp:6:1: remark: Transformed generic-mode kernel to SPMD-mode. [OMP120]
+#pragma omp target teams
+^
+```
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading analysis remark.
+
diff --git a/openmp/docs/remarks/OMP130.md b/openmp/docs/remarks/OMP130.md
index 06872816f664a7..816168995019df 100644
--- a/openmp/docs/remarks/OMP130.md
+++ b/openmp/docs/remarks/OMP130.md
@@ -1,7 +1,6 @@
-.. _omp130:
+(omp130)=
-Removing unused state machine from generic-mode kernel. [OMP130]
-================================================================
+# Removing unused state machine from generic-mode kernel. [OMP130]
This optimization remark indicates that an unused state machine was removed from
a target region. This occurs when there are no parallel regions inside of a
@@ -9,28 +8,27 @@ target construct. Normally, a state machine is required to schedule the threads
inside of a parallel region. If there are no parallel regions, the state machine
is unnecessary because there is only a single thread active at any time.
-Examples
---------
+## Examples
This optimization should occur on any target region that does not contain any
parallel work.
-.. code-block:: c++
+```c++
+void copy(int N, double *X, double *Y) {
+#pragma omp target teams distribute map(tofrom: X[0:N]) map(tofrom: Y[0:N])
+ for (int i = 0; i < N; ++i)
+ Y[i] = X[i];
+}
+```
- void copy(int N, double *X, double *Y) {
- #pragma omp target teams distribute map(tofrom: X[0:N]) map(tofrom: Y[0:N])
- for (int i = 0; i < N; ++i)
- Y[i] = X[i];
- }
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp130.cpp
+omp130.cpp:2:1: remark: Removing unused state machine from generic-mode kernel. [OMP130]
+#pragma omp target teams distribute map(tofrom: X[0:N]) map(tofrom: Y[0:N])
+^
+```
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp130.cpp
- omp130.cpp:2:1: remark: Removing unused state machine from generic-mode kernel. [OMP130]
- #pragma omp target teams distribute map(tofrom: X[0:N]) map(tofrom: Y[0:N])
- ^
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading optimization remark.
+
diff --git a/openmp/docs/remarks/OMP131.md b/openmp/docs/remarks/OMP131.md
index 8b7b999d853699..46a5c743320179 100644
--- a/openmp/docs/remarks/OMP131.md
+++ b/openmp/docs/remarks/OMP131.md
@@ -1,7 +1,6 @@
-Rewriting generic-mode kernel with a customized state machine. [OMP131]
-=======================================================================
+# Rewriting generic-mode kernel with a customized state machine. [OMP131]
-.. _omp131:
+(omp131)=
This optimization remark indicates that a generic-mode kernel on the device was
specialized for the given target region. When offloading in generic-mode, a
@@ -11,38 +10,37 @@ a known number of parallel regions inside the kernel. A much simpler state
machine can be used if it is known that there is no nested parallelism and the
number of regions to schedule is a static amount.
-Examples
---------
+## Examples
This optimization should occur on any generic-mode kernel that has visibility on
all parallel regions, but cannot be moved to SPMD-mode and has no nested
parallelism.
-.. code-block:: c++
-
- #pragma omp declare target
- int TID;
- #pragma omp end declare target
-
- void foo() {
- #pragma omp target
- {
- TID = omp_get_thread_num();
- #pragma omp parallel
- {
- work();
- }
- }
- }
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp131.cpp
- omp131.cpp:8:1: remark: Rewriting generic-mode kernel with a customized state machine. [OMP131]
- #pragma omp target
- ^
-
-Diagnostic Scope
-----------------
+```c++
+#pragma omp declare target
+int TID;
+#pragma omp end declare target
+
+void foo() {
+#pragma omp target
+{
+ TID = omp_get_thread_num();
+ #pragma omp parallel
+ {
+ work();
+ }
+}
+}
+```
+
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp131.cpp
+omp131.cpp:8:1: remark: Rewriting generic-mode kernel with a customized state machine. [OMP131]
+#pragma omp target
+^
+```
+
+## Diagnostic Scope
OpenMP target offloading optimization remark.
+
diff --git a/openmp/docs/remarks/OMP132.md b/openmp/docs/remarks/OMP132.md
index 405c3e64a0e257..674067452ee6a5 100644
--- a/openmp/docs/remarks/OMP132.md
+++ b/openmp/docs/remarks/OMP132.md
@@ -1,7 +1,6 @@
-Generic-mode kernel is executed with a customized state machine that requires a fallback. [OMP132]
-==================================================================================================
+# Generic-mode kernel is executed with a customized state machine that requires a fallback. [OMP132]
-.. _omp132:
+(omp132)=
This analysis remark indicates that a state machine rewrite occurred, but
could not be done fully because of unknown calls to functions that may contain
@@ -10,36 +9,35 @@ worker threads on the device when operating in generic-mode. If there are
unknown parallel regions it prevents the optimization from fully rewriting the
state machine.
-Examples
---------
+## Examples
This will occur for any generic-mode kernel that may contain unknown parallel
-regions. This is typically coupled with the :ref:`OMP133 <omp133>` remark.
-
-.. code-block:: c++
-
- extern void setup();
-
- void foo() {
- #pragma omp target
- {
- setup();
- #pragma omp parallel
- {
- work();
- }
- }
- }
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-analysis=openmp-opt omp132.cpp
- omp133.cpp:4:1: remark: Generic-mode kernel is executed with a customized state machine
- that requires a fallback. [OMP132]
- #pragma omp target
- ^
-
-Diagnostic Scope
-----------------
+regions. This is typically coupled with the {ref}`OMP133 <omp133>` remark.
+
+```c++
+extern void setup();
+
+void foo() {
+#pragma omp target
+{
+ setup();
+ #pragma omp parallel
+ {
+ work();
+ }
+}
+}
+```
+
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-analysis=openmp-opt omp132.cpp
+omp133.cpp:4:1: remark: Generic-mode kernel is executed with a customized state machine
+that requires a fallback. [OMP132]
+#pragma omp target
+^
+```
+
+## Diagnostic Scope
OpenMP target offloading analysis remark.
+
diff --git a/openmp/docs/remarks/OMP133.md b/openmp/docs/remarks/OMP133.md
index 5a734479d495f2..59719dff747a68 100644
--- a/openmp/docs/remarks/OMP133.md
+++ b/openmp/docs/remarks/OMP133.md
@@ -1,70 +1,68 @@
-Call may contain unknown parallel regions. Use `[[omp::assume("omp_no_parallelism")]]` to override. [OMP133]
-====================================================================================================================
+# Call may contain unknown parallel regions. Use `[[omp::assume("omp_no_parallelism")]]` to override. [OMP133]
-.. _omp133:
+(omp133)=
-This analysis remark identifies calls that prevented :ref:`OMP131 <omp131>` from
+This analysis remark identifies calls that prevented {ref}`OMP131 <omp131>` from
providing the generic-mode kernel with a fully specialized state machine. This
remark will identify each call that may contain unknown parallel regions that
caused the kernel to require a fallback.
-Examples
---------
+## Examples
This will occur for any generic-mode kernel that may contain unknown parallel
-regions. This is typically coupled with the :ref:`OMP132 <omp132>` remark.
-
-.. code-block:: c++
-
- extern void setup();
-
- void foo() {
- #pragma omp target
- {
- setup();
- #pragma omp parallel
- {
- work();
- }
- }
- }
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-analysis=openmp-opt omp133.cpp
- omp133.cpp:6:5: remark: Call may contain unknown parallel regions. Use
- `[[omp::assume("omp_no_parallelism")]]` to override. [OMP133]
- setup();
- ^
+regions. This is typically coupled with the {ref}`OMP132 <omp132>` remark.
+
+```c++
+extern void setup();
+
+void foo() {
+#pragma omp target
+{
+ setup();
+ #pragma omp parallel
+ {
+ work();
+ }
+}
+}
+```
+
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass-analysis=openmp-opt omp133.cpp
+omp133.cpp:6:5: remark: Call may contain unknown parallel regions. Use
+`[[omp::assume("omp_no_parallelism")]]` to override. [OMP133]
+setup();
+^
+```
The remark suggests marking the function with the assumption that it contains no
parallel regions. If this is done then the kernel will be rewritten with a fully
specialized state machine.
-.. code-block:: c++
-
- [[omp::assume("omp_no_parallelism")]] extern void setup();
+```c++
+[[omp::assume("omp_no_parallelism")]] extern void setup();
- void foo() {
- #pragma omp target
- {
- setup();
- #pragma omp parallel
- {
- work();
- }
- }
- }
+void foo() {
+#pragma omp target
+{
+ setup();
+ #pragma omp parallel
+ {
+ work();
+ }
+}
+}
+```
-.. code-block:: console
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp133.cpp
+omp133.cpp:4:1: remark: Rewriting generic-mode kernel with a customized state machine. [OMP131]
+#pragma omp target
+^
+```
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O2 -Rpass=openmp-opt omp133.cpp
- omp133.cpp:4:1: remark: Rewriting generic-mode kernel with a customized state machine. [OMP131]
- #pragma omp target
- ^
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP target offloading analysis remark.
+
diff --git a/openmp/docs/remarks/OMP140.md b/openmp/docs/remarks/OMP140.md
index f9bdf257d1a464..4b489c8542f90b 100644
--- a/openmp/docs/remarks/OMP140.md
+++ b/openmp/docs/remarks/OMP140.md
@@ -1,7 +1,6 @@
-.. _omp140:
+(omp140)=
-Could not internalize function. Some optimizations may not be possible. [OMP140]
-====================================================================================================================
+# Could not internalize function. Some optimizations may not be possible. [OMP140]
This analysis remark indicates that function internalization failed for the
given function. Internalization occurs when a call to a function that ordinarily
@@ -10,40 +9,39 @@ only internal visibility. This allows the compiler to make strong static
assertions about the context a function is called in. Without internalization
this analysis would always be invalidated by the possibility of someone calling
the function in a different context outside of the current translation unit.
-This is necessary for optimizations like :ref:`OMP111 <omp111>` and :ref:`OMP120
+This is necessary for optimizations like {ref}`OMP111 <omp111>` and {ref}`OMP120
<omp120>`. If a function failed to be internalized it most likely has linkage
that cannot be copied. Internalization is currently only enabled by default for
OpenMP target offloading.
-Examples
---------
+## Examples
This will occur for any function declaration that has incompatible linkage.
-.. code-block:: c++
-
- __attribute__((weak)) void setup();
-
- void foo() {
- #pragma omp target
- {
- setup();
- #pragma omp parallel
- {
- work();
- }
- }
- }
-
-.. code-block:: console
-
- $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-analysis=openmp-opt omp140.cpp
- omp140.cpp:1:1: remark: Could not internalize function. Some optimizations may not
- be possible. [OMP140]
- __attribute__((weak)) void setup() {
- ^
-
-Diagnostic Scope
-----------------
+```c++
+__attribute__((weak)) void setup();
+
+void foo() {
+#pragma omp target
+{
+ setup();
+ #pragma omp parallel
+ {
+ work();
+ }
+}
+}
+```
+
+```console
+$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass-analysis=openmp-opt omp140.cpp
+omp140.cpp:1:1: remark: Could not internalize function. Some optimizations may not
+be possible. [OMP140]
+__attribute__((weak)) void setup() {
+^
+```
+
+## Diagnostic Scope
OpenMP analysis remark.
+
diff --git a/openmp/docs/remarks/OMP150.md b/openmp/docs/remarks/OMP150.md
index e73eaeb0c4961c..ed8fe9b95d5416 100644
--- a/openmp/docs/remarks/OMP150.md
+++ b/openmp/docs/remarks/OMP150.md
@@ -1,7 +1,6 @@
-.. _omp150:
+(omp150)=
-Parallel region merged with parallel region at <location>. [OMP150]
-===================================================================
+# Parallel region merged with parallel region at \<location>. [OMP150]
This optimization remark indicates that a parallel region was merged with others
into a single parallel region. Parallel region merging fuses consecutive
@@ -10,33 +9,32 @@ the scope of possible OpenMP-specific optimizations within merged parallel
regions. This optimization can also guard sequential code between two parallel
regions if applicable.
-Example
--------
+## Example
This optimization should apply to any compatible and consecutive parallel
regions. In this case the sequential region between the parallel regions will be
guarded so it is only executed by a single thread in the new merged region.
-.. code-block:: c++
+```c++
+void foo() {
+#pragma omp parallel
+ parallel_work();
- void foo() {
- #pragma omp parallel
- parallel_work();
+ sequential_work();
- sequential_work();
+#pragma omp parallel
+ parallel_work();
+}
+```
- #pragma omp parallel
- parallel_work();
- }
+```console
+$ clang++ -fopenmp -O2 -Rpass=openmp-opt -mllvm -openmp-opt-enable-merging omp150.cpp
+omp150.cpp:2:1: remark: Parallel region merged with parallel region at merge.cpp:7:1. [OMP150]
+#pragma omp parallel
+^
+```
-.. code-block:: console
-
- $ clang++ -fopenmp -O2 -Rpass=openmp-opt -mllvm -openmp-opt-enable-merging omp150.cpp
- omp150.cpp:2:1: remark: Parallel region merged with parallel region at merge.cpp:7:1. [OMP150]
- #pragma omp parallel
- ^
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP optimization remark.
+
diff --git a/openmp/docs/remarks/OMP160.md b/openmp/docs/remarks/OMP160.md
index 5012910b18cc31..f0a63ef96417ec 100644
--- a/openmp/docs/remarks/OMP160.md
+++ b/openmp/docs/remarks/OMP160.md
@@ -1,7 +1,6 @@
-.. _omp160:
+(omp160)=
-Removing parallel region with no side-effects. [OMP160]
-=======================================================
+# Removing parallel region with no side-effects. [OMP160]
This optimization remark indicates that a parallel region was deleted because it
was not found to have any side-effects. This can occur if the region does not
@@ -10,35 +9,34 @@ is necessary because the barrier between sequential and parallel code typically
prevents dead code elimination from completely removing the region. Otherwise
there will still be overhead to fork and merge the threads with no work done.
-Example
--------
+## Example
This optimization occurs whenever a parallel region was not found to have any
side-effects. This can occur if the parallel region only reads memory or is
simply empty.
-.. code-block:: c++
-
- void foo() {
- #pragma omp parallel
- { }
- #pragma omp parallel
- { int x = 1; }
- }
- }
-
-.. code-block:: console
-
- $ clang++ -fopenmp -O2 -Rpass=openmp-opt omp160.cpp
- omp160.cpp:4:1: remark: Removing parallel region with no side-effects. [OMP160] [-Rpass=openmp-opt]
- #pragma omp parallel
- ^
- delete.cpp:2:1: remark: Removing parallel region with no side-effects. [OMP160] [-Rpass=openmp-opt]
- #pragma omp parallel
- ^
- ^
-
-Diagnostic Scope
-----------------
+```c++
+void foo() {
+#pragma omp parallel
+ { }
+#pragma omp parallel
+ { int x = 1; }
+}
+}
+```
+
+```console
+$ clang++ -fopenmp -O2 -Rpass=openmp-opt omp160.cpp
+omp160.cpp:4:1: remark: Removing parallel region with no side-effects. [OMP160] [-Rpass=openmp-opt]
+#pragma omp parallel
+^
+delete.cpp:2:1: remark: Removing parallel region with no side-effects. [OMP160] [-Rpass=openmp-opt]
+#pragma omp parallel
+^
+^
+```
+
+## Diagnostic Scope
OpenMP optimization remark.
+
diff --git a/openmp/docs/remarks/OMP170.md b/openmp/docs/remarks/OMP170.md
index f33c7a6921dfc1..7312a3f71a7a8b 100644
--- a/openmp/docs/remarks/OMP170.md
+++ b/openmp/docs/remarks/OMP170.md
@@ -1,7 +1,6 @@
-.. _omp170:
+(omp170)=
-OpenMP runtime call <call> deduplicated. [OMP170]
-====================================================================
+# OpenMP runtime call \<call> deduplicated. [OMP170]
This optimization remark indicates that a call to an OpenMP runtime call was
replaced with the result of an existing one. This occurs when the compiler knows
@@ -10,32 +9,31 @@ by replacing all calls to that function with the result of the first call. This
cannot be done automatically by the compiler because the implementations of the
OpenMP runtime calls live in a separate library the compiler cannot see.
-Example
--------
+## Example
This optimization will trigger for known OpenMP runtime calls whose return value
will not change.
-.. code-block:: c++
-
- void foo(int N) {
- double *A = malloc(N * omp_get_thread_limit());
- double *B = malloc(N * omp_get_thread_limit());
+```c++
+void foo(int N) {
+ double *A = malloc(N * omp_get_thread_limit());
+ double *B = malloc(N * omp_get_thread_limit());
- #pragma omp parallel
- work(&A[omp_get_thread_num() * N]);
- #pragma omp parallel
- work(&B[omp_get_thread_num() * N]);
- }
+#pragma omp parallel
+ work(&A[omp_get_thread_num() * N]);
+#pragma omp parallel
+ work(&B[omp_get_thread_num() * N]);
+}
+```
-.. code-block:: console
+```console
+$ clang -fopenmp -O2 -Rpass=openmp-opt omp170.c
+ompi170.c:2:26: remark: OpenMP runtime call omp_get_thread_limit deduplicated. [OMP170]
+double *A = malloc(N * omp_get_thread_limit());
+ ^
+```
- $ clang -fopenmp -O2 -Rpass=openmp-opt omp170.c
- ompi170.c:2:26: remark: OpenMP runtime call omp_get_thread_limit deduplicated. [OMP170]
- double *A = malloc(N * omp_get_thread_limit());
- ^
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP optimization remark.
+
diff --git a/openmp/docs/remarks/OMP180.md b/openmp/docs/remarks/OMP180.md
index c2d937236b324c..36513d1db57de0 100644
--- a/openmp/docs/remarks/OMP180.md
+++ b/openmp/docs/remarks/OMP180.md
@@ -1,38 +1,36 @@
-.. _omp180:
+(omp180)=
-Replacing OpenMP runtime call <call> with <value>.
-====================================================================
+# Replacing OpenMP runtime call \<call> with \<value>.
This optimization remark indicates that analysis determined an OpenMP runtime
calls can be replaced with a constant value. This can occur when an OpenMP
runtime call that queried some internal state was found to always return a
single value after analysis.
-Example
--------
+## Example
This optimization will trigger for most target regions to simplify the runtime
once certain constants are known. This will trigger for internal runtime
functions so it requires enabling verbose remarks with
`-openmp-opt-verbose-remarks` (prefixed with `-mllvm` for use with clang).
-.. code-block:: c++
+```c++
+void foo() {
+#pragma omp target parallel
+ { }
+}
+```
- void foo() {
- #pragma omp target parallel
- { }
- }
+```console
+$ clang test.c -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass=openmp-opt \
+ -mllvm -openmp-opt-verbose-remarks
+remark: Replacing runtime call __kmpc_is_spmd_exec_mode with 1. [OMP180] [-Rpass=openmp-opt]
+remark: Replacing runtime call __kmpc_is_spmd_exec_mode with 1. [OMP180] [-Rpass=openmp-opt]
+remark: Replacing runtime call __kmpc_parallel_level with 1. [OMP180] [-Rpass=openmp-opt]
+remark: Replacing runtime call __kmpc_parallel_level with 1. [OMP180] [-Rpass=openmp-opt]
+```
-.. code-block:: console
-
- $ clang test.c -fopenmp -fopenmp-targets=nvptx64 -O1 -Rpass=openmp-opt \
- -mllvm -openmp-opt-verbose-remarks
- remark: Replacing runtime call __kmpc_is_spmd_exec_mode with 1. [OMP180] [-Rpass=openmp-opt]
- remark: Replacing runtime call __kmpc_is_spmd_exec_mode with 1. [OMP180] [-Rpass=openmp-opt]
- remark: Replacing runtime call __kmpc_parallel_level with 1. [OMP180] [-Rpass=openmp-opt]
- remark: Replacing runtime call __kmpc_parallel_level with 1. [OMP180] [-Rpass=openmp-opt]
-
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP optimization remark.
+
diff --git a/openmp/docs/remarks/OMP190.md b/openmp/docs/remarks/OMP190.md
index 2170c0acf97580..cbcd9ebbeb5d02 100644
--- a/openmp/docs/remarks/OMP190.md
+++ b/openmp/docs/remarks/OMP190.md
@@ -1,7 +1,6 @@
-.. _omp190:
+(omp190)=
-Redundant barrier eliminated. (device only)
-====================================================================
+# Redundant barrier eliminated. (device only)
This optimization remark indicates that analysis determined an aligned
barrier in the device code to be redundant. This can occur when state
@@ -9,15 +8,14 @@ updates that have been synchronized by the barrier were eliminated too.
See also "Co-Designing an OpenMP GPU Runtime and Optimizations for Near-Zero
Overhead Execution", IPDPS'22.
-Example
--------
+## Example
This optimization will trigger for most target regions if state initialization
was removed as a consequence of "state forwarding". This will trigger for
internal runtime functions so it requires enabling verbose remarks with
`-openmp-opt-verbose-remarks` (prefixed with `-mllvm` for use with clang).
-Diagnostic Scope
-----------------
+## Diagnostic Scope
OpenMP optimization remark.
+
diff --git a/openmp/docs/remarks/OptimizationRemarks.md b/openmp/docs/remarks/OptimizationRemarks.md
index 2c683a4376c491..eb14376833c1cd 100644
--- a/openmp/docs/remarks/OptimizationRemarks.md
+++ b/openmp/docs/remarks/OptimizationRemarks.md
@@ -1,47 +1,44 @@
-OpenMP Optimization Remarks
-===========================
+# OpenMP Optimization Remarks
-The :doc:`OpenMP-Aware optimization pass </optimizations/OpenMPOpt>` is able to
+The {doc}`OpenMP-Aware optimization pass </optimizations/OpenMPOpt>` is able to
generate compiler remarks for performed and missed optimisations. To emit them,
-pass these options to the Clang invocation: ``-Rpass=openmp-opt
--Rpass-analysis=openmp-opt -Rpass-missed=openmp-opt``. For more information and
+pass these options to the Clang invocation: `-Rpass=openmp-opt
+-Rpass-analysis=openmp-opt -Rpass-missed=openmp-opt`. For more information and
features of the remark system, consult the clang documentation:
-+ `Clang options to emit optimization reports <https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports>`_
-+ `Clang diagnostic and remark flags <https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags>`_
-+ The `-foptimization-record-file flag
- <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file>`_
- and the `-fsave-optimization-record flag
- <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record>`_
+- [Clang options to emit optimization reports](https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports)
+- [Clang diagnostic and remark flags](https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags)
+- The [-foptimization-record-file flag](https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file)
+ and the [-fsave-optimization-record flag](https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record)
+## OpenMP Remarks
-OpenMP Remarks
---------------
+```{toctree}
+:hidden: true
+:maxdepth: 1
-.. toctree::
- :hidden:
- :maxdepth: 1
-
- OMP100
- OMP101
- OMP102
- OMP110
- OMP111
- OMP112
- OMP113
- OMP120
- OMP121
- OMP130
- OMP131
- OMP132
- OMP133
- OMP140
- OMP150
- OMP160
- OMP170
- OMP180
- OMP190
+OMP100
+OMP101
+OMP102
+OMP110
+OMP111
+OMP112
+OMP113
+OMP120
+OMP121
+OMP130
+OMP131
+OMP132
+OMP133
+OMP140
+OMP150
+OMP160
+OMP170
+OMP180
+OMP190
+```
+```{eval-rst}
.. list-table::
:widths: 15 15 70
:header-rows: 1
@@ -115,3 +112,5 @@ OpenMP Remarks
* - :ref:`OMP190 <omp190>`
- Optimization
- Redundant barrier eliminated. (device only)
+```
+
>From 50e5f59cf8eab7683f27521b3c0fc5e8269a04f6 Mon Sep 17 00:00:00 2001
From: Reid Kleckner <rkleckner at nvidia.com>
Date: Sat, 5 Sep 2026 04:45:33 +0000
Subject: [PATCH 2/3] [OpenMP][docs] Finish MyST migration for OpenMP docs
---
openmp/docs/CommandLineArgumentReference.md | 203 +++++++-------
openmp/docs/README.txt | 18 +-
openmp/docs/SupportAndFAQ.md | 13 +-
openmp/docs/design/Offloading.md | 4 +-
openmp/docs/design/Overview.md | 5 +-
openmp/docs/design/Runtimes.md | 288 +++++---------------
openmp/docs/index.md | 55 ++--
openmp/docs/openacc/OpenMPExtensions.md | 7 +-
openmp/docs/openacc/Overview.md | 5 +-
openmp/docs/optimizations/OpenMPOpt.md | 12 +-
openmp/docs/optimizations/Overview.md | 13 +-
openmp/docs/remarks/OMP100.md | 9 +-
openmp/docs/remarks/OMP113.md | 2 +-
openmp/docs/remarks/OMP121.md | 3 +-
openmp/docs/remarks/OMP133.md | 3 +-
openmp/docs/remarks/OMP150.md | 1 -
openmp/docs/remarks/OMP180.md | 3 +-
openmp/docs/remarks/OMP190.md | 2 +-
openmp/docs/remarks/OptimizationRemarks.md | 64 +++--
19 files changed, 258 insertions(+), 452 deletions(-)
diff --git a/openmp/docs/CommandLineArgumentReference.md b/openmp/docs/CommandLineArgumentReference.md
index 3b91d4355464d0..40476478817448 100644
--- a/openmp/docs/CommandLineArgumentReference.md
+++ b/openmp/docs/CommandLineArgumentReference.md
@@ -1,3 +1,6 @@
+:::{program} clang
+:::
+
# OpenMP Command-Line Argument Reference
Welcome to the OpenMP in LLVM command line argument reference. The content is
@@ -7,171 +10,155 @@ Section {ref}`general_command_line_arguments` lists OpenMP command line options
for multicore programming while {ref}`offload_command_line_arguments` lists
options relevant to OpenMP target offloading.
-(general-command-line-arguments)=
+(general_command_line_arguments)=
## OpenMP Command-Line Arguments
-### `-fopenmp`
+(fopenmp)=
+:::{option} -fopenmp
Enable the OpenMP compilation toolchain. The compiler will parse OpenMP
compiler directives and generate parallel code.
+:::
-### `-fopenmp-extensions`
+(fopenmp-extensions)=
+:::{option} -fopenmp-extensions
Enable all `Clang` extensions for OpenMP directives and clauses. A list of
current extensions and their implementation status can be found on the
[support](https://clang.llvm.org/docs/OpenMPSupport.html#openmp-extensions)
page.
+:::
-### `-fopenmp-simd`
+(fopenmp-simd)=
+:::{option} -fopenmp-simd
This option enables OpenMP only for single instruction, multiple data
(SIMD) constructs.
+:::
-### `-static-openmp`
+(static-openmp)=
+:::{option} -static-openmp
Use the static OpenMP host runtime while linking.
+:::
-### `-fopenmp-version=<arg>`
+(fopenmp-version-arg)=
+:::{option} -fopenmp-version=<arg>
Set the OpenMP version to a specific version `<arg>` of the OpenMP standard.
For example, you may use `-fopenmp-version=45` to select version 4.5 of
the OpenMP standard. The default value is `-fopenmp-version=51` for `Clang`.
+:::
-(offload-command-line-arguments)=
+(offload_command_line_arguments)=
## Offloading Specific Command-Line Arguments
(fopenmp-targets)=
-### `-fopenmp-targets`
-
+:::{option} -fopenmp-targets
Specify which OpenMP offloading targets should be supported. For example, you
-may specify
-
-`-fopenmp-targets=amdgcn-amd-amdhsa,nvptx64`
-
-. This option is
-often optional when
-
-{ref}`offload_arch`
-
- is provided.
+may specify `-fopenmp-targets=amdgcn-amd-amdhsa,nvptx64`. This option is
+often optional when {ref}`--offload-arch <offload_arch>` is provided.
It is also possible to offload to CPU architectures, for instance with
+`-fopenmp-targets=x86_64-pc-linux-gnu`.
+:::
+(offload_arch)=
-`-fopenmp-targets=x86_64-pc-linux-gnu`
-
-.
-
-(offload-arch)=
-
-### `--offload-arch`
-
+:::{option} --offload-arch
Specify the device architecture for OpenMP offloading. For instance
+`--offload-arch=sm_80` to target an Nvidia Tesla A100,
+`--offload-arch=gfx90a` to target an AMD Instinct MI250X, or
+`--offload-arch=sm_80,gfx90a` to target both.
-
-`--offload-arch=sm_80`
-
- to target an Nvidia Tesla A100,
-
-
-`--offload-arch=gfx90a`
-
- to target an AMD Instinct MI250X, or
-
-
-`--offload-arch=sm_80,gfx90a`
-
- to target both.
-
-It is also possible to specify
-
-{ref}`fopenmp-targets`
-
- without specifying
-
-
-`--offload-arch`
-
-. In that case, the executables
-
-`amdgpu-arch`
-
- or
-
-
-`nvptx-arch`
-
- will be executed as part of the compiler driver to
+It is also possible to specify {option}`-fopenmp-targets` without specifying
+`--offload-arch`. In that case, the executables `amdgpu-arch` or
+`nvptx-arch` will be executed as part of the compiler driver to
detect the device architecture automatically.
Finally, the device architecture will also be automatically inferred with
+`--offload-arch=native`.
+:::
+(offload-device-only)=
-`--offload-arch=native`
-
-.
-
-### `--offload-device-only`
-
+:::{option} --offload-device-only
Compile only the code that goes on the device. This option is mainly for
debugging purposes. It is primarily used for inspecting the intermediate
representation (IR) output when compiling for the device. It may also be used
if device-only runtimes are created.
+:::
-### `--offload-host-only`
+(offload-host-only)=
+:::{option} --offload-host-only
Compile only the code that goes on the host. With this option enabled, the
`.llvm.offloading` section with embedded device code will not be included in
the intermediate representation.
+:::
-### `--offload-host-device`
+(offload-host-device)=
+:::{option} --offload-host-device
Compile the target regions for both the host and the device. That is the
default option.
+:::
-### `-Xopenmp-target <arg>`
+(xopenmp-target-arg)=
+(xopenmp-target-triple-arg)=
+
+:::{option} -Xopenmp-target[=<triple>] <arg>
Pass an argument `<arg>` to the offloading toolchain, for instance
`-Xopenmp-target -march=sm_80`.
-### `-Xopenmp-target=<triple> <arg>`
-
Pass an argument `<arg>` to the offloading toolchain for the target
`<triple>`. That is especially useful when an argument must differ for each
triple. For instance `-Xopenmp-target=nvptx64 --offload-arch=sm_80
-Xopenmp-target=amdgcn --offload-arch=gfx90a` to specify the device
-architecture. Alternatively, {ref}`Xarch_host` and {ref}`Xarch_device` can
+architecture. Alternatively, {ref}`-Xarch_host <Xarch_host>` and
+{ref}`-Xarch_device <Xarch_device>` can
pass an argument to the host and device compilation toolchain.
+:::
-### `-Xoffload-linker<triple> <arg>`
+(xoffload-linker-triple-arg)=
+:::{option} -Xoffload-linker<triple> <arg>
Pass an argument `<arg>` to the offloading linker for the target specified in
`<triple>`.
+:::
-(xarch-device)=
+(Xarch_device)=
-### `-Xarch_device <arg>`
+(xarch-device-arg)=
+:::{option} -Xarch_device <arg>
Pass an argument `<arg>` to the device compilation toolchain.
+:::
-(xarch-host)=
+(Xarch_host)=
-### `-Xarch_host <arg>`
+(xarch-host-arg)=
+:::{option} -Xarch_host <arg>
Pass an argument `<arg>` to the host compilation toolchain.
+:::
-### `-foffload-lto[=<arg>]`
+(foffload-lto-arg)=
+:::{option} -foffload-lto[=<arg>]
Enable device link time optimization (LTO) and select the LTO mode `<arg>`.
Select either `-foffload-lto=thin` or `-foffload-lto=full`. Thin LTO takes
less time while still achieving some performance gains. If no argument is set,
this option defaults to `-foffload-lto=full`.
+:::
-### `-fopenmp-offload-mandatory`
+(fopenmp-offload-mandatory)=
+:::{option} -fopenmp-offload-mandatory
This option is set to avoid generating the host fallback code
executed when offloading to the device fails. That is
helpful when the target contains code that cannot be compiled for the host, for
@@ -181,15 +168,13 @@ This option can also be used to reduce compile time.
This option should not be used when one wants to verify that the code is being
offloaded to the device. Instead, set the environment variable
-
-
-`OMP_TARGET_OFFLOAD='MANDATORY'`
-
- to confirm that the code is being offloaded to
+`OMP_TARGET_OFFLOAD='MANDATORY'` to confirm that the code is being offloaded to
the device.
+:::
-### `-fopenmp-target-debug[=<arg>]`
+(fopenmp-target-debug-arg)=
+:::{option} -fopenmp-target-debug[=<arg>]
Enable debugging in the device runtime library (RTL). Note that it is both
necessary to configure the debugging in the device runtime at compile-time with
`-fopenmp-target-debug=<arg>` and enable debugging at runtime with the
@@ -199,57 +184,51 @@ environment variable `LIBOMPTARGET_DEBUG` can be set to debug both Nvidia and
AMD GPU targets. For more information, see the
[debugging instructions](https://openmp.llvm.org/design/Runtimes.html#debugging).
The debugging instructions list the supported debugging arguments.
+:::
-### `-fopenmp-target-jit`
+(fopenmp-target-jit)=
+:::{option} -fopenmp-target-jit
Emit code that is Just-in-Time (JIT) compiled for OpenMP offloading. Embed
LLVM-IR for the device code in the object files rather than binary code for the
respective target. At runtime, the LLVM-IR is optimized again and compiled for
the target device. The optimization level can be set at runtime with
-
-
-`LIBOMPTARGET_JIT_OPT_LEVEL`
-
-, for instance,
-
-
-`LIBOMPTARGET_JIT_OPT_LEVEL=3`
-
- corresponding to optimizations level
-
-`-O3`
-
-.
+`LIBOMPTARGET_JIT_OPT_LEVEL`, for instance,
+`LIBOMPTARGET_JIT_OPT_LEVEL=3` corresponding to optimizations level `-O3`.
See the
-
-
[OpenMP JIT details](https://openmp.llvm.org/design/Runtimes.html#libomptarget-jit-pre-opt-ir-module)
-
-
for instructions on extracting the embedded device code before or after the
JIT and more.
We want to emphasize that JIT for OpenMP offloading is good for debugging as
the target IR can be extracted, modified, and injected at runtime.
+:::
-### `--offload-new-driver`
+(offload-new-driver)=
+:::{option} --offload-new-driver
Deprecated and accepted as a no-op. The new offloading driver is always used for
offloading compilation; the legacy driver has been removed.
+:::
-### `--offload-link`
+(offload-link)=
-Use the new offloading linker `clang-linker-wrapper` to perform the link job.
-`clang-linker-wrapper` is the default offloading linker for OpenMP. This option
+:::{option} --offload-link
+Use the new offloading linker {title-reference}`clang-linker-wrapper` to perform the link job.
+{title-reference}`clang-linker-wrapper` is the default offloading linker for OpenMP. This option
can be used to use the new offloading linker in toolchains that do not automatically
use it. It is necessary to enable this option when linking with CUDA or HIP files.
+:::
-### `-nogpulib`
+(nogpulib)=
+:::{option} -nogpulib
Do not link the device library for CUDA or HIP device compilation.
+:::
-### `-nogpuinc`
+(nogpuinc)=
+:::{option} -nogpuinc
Do not include the default CUDA or HIP headers, and do not add CUDA or HIP
include paths.
-
+:::
diff --git a/openmp/docs/README.txt b/openmp/docs/README.txt
index fe53bcec03ba40..27d478d001d68c 100644
--- a/openmp/docs/README.txt
+++ b/openmp/docs/README.txt
@@ -1,9 +1,9 @@
OpenMP LLVM Documentation
==================
-OpenMP LLVM's documentation is written in reStructuredText, a lightweight
-plaintext markup language (file extension `.rst`). While the
-reStructuredText documentation should be quite readable in source form, it
+OpenMP LLVM's documentation is written in MyST Markdown, a lightweight
+plaintext markup language (file extension `.md`). While the
+Markdown documentation should be quite readable in source form, it
is mostly meant to be processed by the Sphinx documentation generation
system to create HTML pages which are hosted on <https://llvm.org/docs/> and
updated after every commit. Manpage output is also supported, see below.
@@ -17,13 +17,13 @@ Sphinx <http://sphinx-doc.org/> and then do:
$BROWSER <build-dir>/docs/html/index.html
The mapping between reStructuredText files and generated documentation is
-`docs/Foo.rst` <-> `<build-dir>/projects/openmp/docs//html/Foo.html` <->
+`docs/Foo.md` <-> `<build-dir>/projects/openmp/docs//html/Foo.html` <->
`https://openmp.llvm.org/docs/Foo.html`.
If you are interested in writing new documentation, you will want to read
`llvm/docs/SphinxQuickstartTemplate.md` which will get you writing
-documentation very fast and includes examples of the most important
-reStructuredText markup syntax.
+documentation very fast and includes examples of the most important MyST
+Markdown syntax.
Manpage Output
===============
@@ -38,8 +38,8 @@ directory `<build-dir>/docs/man/`.
make
man -l >build-dir>/docs/man/FileCheck.1
-The correspondence between .rst files and man pages is
-`docs/CommandGuide/Foo.rst` <-> `<build-dir>/projects/openmp/docs//man/Foo.1`.
-These .rst files are also included during HTML generation so they are also
+The correspondence between .md files and man pages is
+`docs/CommandGuide/Foo.md` <-> `<build-dir>/projects/openmp/docs//man/Foo.1`.
+These .md files are also included during HTML generation so they are also
viewable online (as noted above) at e.g.
`https://openmp.llvm.org/docs/CommandGuide/Foo.html`.
diff --git a/openmp/docs/SupportAndFAQ.md b/openmp/docs/SupportAndFAQ.md
index ee25f9566faee4..02d4bcd2758882 100644
--- a/openmp/docs/SupportAndFAQ.md
+++ b/openmp/docs/SupportAndFAQ.md
@@ -112,15 +112,15 @@ For now, the answer is most likely *no*. Please see {ref}`build_offload_capable_
For now, the answer is most likely *no*. Please see {ref}`build_offload_capable_compiler`.
-(math-and-complex-in-target-regions)=
+(math_and_complex_in_target_regions)=
-### Q: Does Clang support `<math.h>` and `<complex.h>` operations in OpenMP target on GPUs?
+### Q: Does Clang support {title-reference}`<math.h>` and {title-reference}`<complex.h>` operations in OpenMP target on GPUs?
Yes, LLVM/Clang allows math functions and complex arithmetic inside of OpenMP
target regions that are compiled for GPUs.
-Clang provides a set of wrapper headers that are found first when `math.h` and
-`complex.h`, for C, `cmath` and `complex`, for C++, or similar headers are
+Clang provides a set of wrapper headers that are found first when {title-reference}`math.h` and
+{title-reference}`complex.h`, for C, {title-reference}`cmath` and {title-reference}`complex`, for C++, or similar headers are
included by the application. These wrappers will eventually include the system
version of the corresponding header file after setting up a target device
specific environment. The fact that the system header is included is important
@@ -128,7 +128,7 @@ because they differ based on the architecture and operating system and may
contain preprocessor, variable, and function definitions that need to be
available in the target region regardless of the targeted device architecture.
However, various functions may require specialized device versions, e.g.,
-`sin`, and others are only available on certain devices, e.g., `__umul64hi`. To
+{title-reference}`sin`, and others are only available on certain devices, e.g., {title-reference}`__umul64hi`. To
provide "native" support for math and complex on the respective architecture,
Clang will wrap the "native" math functions, e.g., as provided by the device
vendor, in an OpenMP begin/end declare variant. These functions will then be
@@ -276,7 +276,7 @@ clang++ openmp.cpp -fopenmp --offload-arch=gfx90a -Xoffload-linker -lc
```
For more information on how this is implemented in LLVM/OpenMP's offloading
-runtime, refer to the [runtime documentation](libomptarget_libc).
+runtime, refer to the {ref}`runtime documentation <libomptarget_libc>`.
### Q: What command line options can I use for OpenMP?
@@ -308,4 +308,3 @@ compilations. Consider enabling `Ccache` with
Feel free to post questions or browse old threads at
[LLVM Discourse](https://discourse.llvm.org/c/runtimes/openmp/).
-
diff --git a/openmp/docs/design/Offloading.md b/openmp/docs/design/Offloading.md
index 85af57e4ebf68f..068a658a1a3833 100644
--- a/openmp/docs/design/Offloading.md
+++ b/openmp/docs/design/Offloading.md
@@ -1,11 +1,11 @@
# OpenMP in LLVM --- Offloading Design
-```{toctree}
+:::{toctree}
:glob: true
:hidden: true
:maxdepth: 1
GPUSPMDMode
GPUGenericMode
-```
+:::
diff --git a/openmp/docs/design/Overview.md b/openmp/docs/design/Overview.md
index 90b0029fda55e3..59f8b90c41b41c 100644
--- a/openmp/docs/design/Overview.md
+++ b/openmp/docs/design/Overview.md
@@ -4,12 +4,11 @@
- OpenMP Booth @ SC19: "OpenMP clang and flang Development" <https://youtu.be/6yOa-hRi63M>
-```{toctree}
+:::{toctree}
:glob: true
:hidden: true
:maxdepth: 1
Runtimes
Offloading
-```
-
+:::
diff --git a/openmp/docs/design/Runtimes.md b/openmp/docs/design/Runtimes.md
index e31dcaf9d5e9cd..6fae6a84a549a6 100644
--- a/openmp/docs/design/Runtimes.md
+++ b/openmp/docs/design/Runtimes.md
@@ -1,4 +1,4 @@
-(openmp-runtimes)=
+(openmp_runtimes)=
# LLVM/OpenMP Runtimes
@@ -17,7 +17,7 @@ For general information on debugging OpenMP target offloading applications, see
An [early (2015) design document](https://raw.githubusercontent.com/llvm/llvm-project/main/openmp/runtime/doc/Reference.pdf)
for the LLVM/OpenMP host runtime, aka. `libomp.so`, is available as a [pdf](https://raw.githubusercontent.com/llvm/llvm-project/main/openmp/runtime/doc/Reference.pdf).
-(libomp-environment-vars)=
+(libomp_environment_vars)=
### Environment Variables
@@ -65,21 +65,13 @@ is used.
Enables (`true`) or disables (`false`) the dynamic adjustment of the
number of threads.
-**Default:**
-
-
-
-`false`
+**Default:** `false`
#### OMP_MAX_ACTIVE_LEVELS
The maximum number of levels of parallel nesting for the program.
-**Default:**
-
-
-
-`1`
+**Default:** `1`
#### OMP_NESTED
@@ -89,11 +81,7 @@ Deprecated. Please use `OMP_MAX_ACTIVE_LEVELS` to control nested parallelism
Enables (`true`) or disables (`false`) nested parallelism.
-**Default:**
-
-
-
-`false`
+**Default:** `false`
#### OMP_NUM_THREADS
@@ -112,21 +100,9 @@ list is left out, it implies the normal default value for threads is used at the
outer-most level. If the integer is left out of any other level, the number of
threads for that level is inherited from the previous level.
-**Default:**
-
- The number of processors visible to the operating system on which the program is executed.
-
-**Syntax:**
-
-
-
-`OMP_NUM_THREADS=value[,value]*`
-
-**Example:**
-
-
-
-`OMP_NUM_THREADS=4,3`
+**Default:** The number of processors visible to the operating system on which the program is executed.\
+**Syntax:** `OMP_NUM_THREADS=value[,value]*`\
+**Example:** `OMP_NUM_THREADS=4,3`
#### OMP_PLACES
@@ -232,35 +208,14 @@ primary thread is bound.
If set to `spread`, the primary thread's partition is subdivided and threads
are bound to single place successive sub-partitions.
-**Related environment variables:**
-
-
-
-`KMP_AFFINITY`
-
- (overrides
-
-`OMP_PROC_BIND`
-
-).
+**Related environment variables:** `KMP_AFFINITY` (overrides `OMP_PROC_BIND`).
#### OMP_SCHEDULE
Sets the run-time schedule type and an optional chunk size.
-**Default:**
-
-
-
-`static`
-
-, no chunk size specified
-
-**Syntax:**
-
-
-
-`OMP_SCHEDULE="kind[,chunk_size]"`
+**Default:** `static`, no chunk size specified\
+**Syntax:** `OMP_SCHEDULE="kind[,chunk_size]"`
#### OMP_STACKSIZE
@@ -285,23 +240,8 @@ has no effect.
- 32-bit architecture: `2M`
- 64-bit architecture: `4M`
-**Related environment variables:**
-
-
-
-`KMP_STACKSIZE`
-
- (overrides
-
-`OMP_STACKSIZE`
-
-).
-
-**Example:**
-
-
-
-`OMP_STACKSIZE=8M`
+**Related environment variables:** `KMP_STACKSIZE` (overrides `OMP_STACKSIZE`).\
+**Example:** `OMP_STACKSIZE=8M`
#### OMP_THREAD_LIMIT
@@ -315,21 +255,8 @@ the team was reduced, but the program will continue.
The `omp_get_thread_limit()` routine returns the value of the limit.
-**Default:**
-
- No enforced limit
-
-**Related environment variable:**
-
-
-
-`KMP_ALL_THREADS`
-
- (overrides
-
-`OMP_THREAD_LIMIT`
-
-).
+**Default:** No enforced limit\
+**Related environment variable:** `KMP_ALL_THREADS` (overrides `OMP_THREAD_LIMIT`).
#### OMP_WAIT_POLICY
@@ -337,11 +264,7 @@ Decides whether threads spin (active) or yield (passive) while they are waiting.
`OMP_WAIT_POLICY=active` is an alias for `KMP_LIBRARY=turnaround`, and
`OMP_WAIT_POLICY=passive` is an alias for `KMP_LIBRARY=throughput`.
-**Default:**
-
-
-
-`passive`
+**Default:** `passive`
:::{note}
Although the default is `passive`, unless the user has explicitly set
@@ -422,23 +345,8 @@ considered a separate level for the sort operations.
The `offset` specifier indicates the starting position for thread assignment.
-**Default:**
-
-
-
-`noverbose,warnings,respect,granularity=core,none`
-
-**Related environment variable:**
-
-
-
-`OMP_PROC_BIND`
-
- (
-
-`KMP_AFFINITY`
-
- takes precedence)
+**Default:** `noverbose,warnings,respect,granularity=core,none`\
+**Related environment variable:** `OMP_PROC_BIND` (`KMP_AFFINITY` takes precedence)
:::{note}
On Windows with multiple processor groups, the norespect affinity modifier
@@ -467,8 +375,8 @@ and any affinity API calls.
The following `modifiers` are ignored in `KMP_HIDDEN_HELPER_AFFINITY` and are only valid
for `KMP_AFFINITY`:
-\* `respect` and `norespect`
-\* `reset` and `noreset`
+* `respect` and `norespect`
+* `reset` and `noreset`
#### KMP_ALL_THREADS
@@ -479,21 +387,8 @@ message. If this limit is reached at the time an OpenMP parallel region begins,
a one-time warning message may be generated indicating that the number of
threads in the team was reduced, but the program will continue execution.
-**Default:**
-
- No enforced limit.
-
-**Related environment variable:**
-
-
-
-`OMP_THREAD_LIMIT`
-
- (
-
-`KMP_ALL_THREADS`
-
- takes precedence)
+**Default:** No enforced limit.\
+**Related environment variable:** `OMP_THREAD_LIMIT` (`KMP_ALL_THREADS` takes precedence)
#### KMP_BLOCKTIME
@@ -505,26 +400,14 @@ specify/change the units. Defaults units is milliseconds.
Specify `infinite` for an unlimited wait time.
-**Default:**
-
- 200 milliseconds
-
-**Related Environment Variable:**
-
-
-
-`KMP_LIBRARY`
-
-**Example:**
-
-
-
-`KMP_BLOCKTIME=1ms`
+**Default:** 200 milliseconds\
+**Related Environment Variable:** `KMP_LIBRARY`\
+**Example:** `KMP_BLOCKTIME=1ms`
#### KMP_CPUINFO_FILE
Specifies an alternate file name for a file containing the machine topology
-description. The file must be in the same format as {file}`/proc/cpuinfo`.
+description. The file must be in the same format as `/proc/cpuinfo`.
**Default:** None
@@ -537,17 +420,8 @@ a given parallel region, for a given data set and reduction operation, a
floating point reduction done for an OpenMP reduction clause has a consistent
floating point result from run to run, since round-off errors are identical.
-**Default:**
-
-
-
-`false`
-
-**Example:**
-
-
-
-`KMP_DETERMINISTIC_REDUCTION=true`
+**Default:** `false`\
+**Example:** `KMP_DETERMINISTIC_REDUCTION=true`
#### KMP_DYNAMIC_MODE
@@ -607,37 +481,17 @@ An optional colon (:) can be specified at the beginning of the syntax to specify
Supported unit IDs are not case-insensitive.
-`S`
-
- - socket
-
-`num_units`
-
- specifies the requested number of sockets.
-
-`D`
-
- - die
-
-`num_units`
-
- specifies the requested number of dies per socket.
-
-`C`
-
- - core
-
-`num_units`
+`S` - socket\
+`num_units` specifies the requested number of sockets.
- specifies the requested number of cores per die - if any - otherwise, per socket.
+`D` - die\
+`num_units` specifies the requested number of dies per socket.
-`T`
+`C` - core\
+`num_units` specifies the requested number of cores per die - if any - otherwise, per socket.
- - thread
-
-`num_units`
-
- specifies the requested number of HW threads per core.
+`T` - thread\
+`num_units` specifies the requested number of HW threads per core.
:::{note}
`num_units` can be left out or explicitly specified as `*` instead of a positive integer
@@ -650,7 +504,7 @@ e.g., `1s,*c` means use 1 socket and all the cores on that socket
`attribute` - (Optional) An attribute differentiating resources at a particular level. The attributes available to users are:
- **Core type** - On Intel architectures, this can be `intel_atom` or `intel_core`
-- **Core efficiency** - This is specified as `eff`{emphasis}`num` where {emphasis}`num` is a number from 0
+- **Core efficiency** - This is specified as `eff`*num* where *num* is a number from 0
to the number of core efficiencies detected in the machine topology minus one.
E.g., `eff0`. The greater the efficiency number the more performant the core. There may be
more core efficiencies than core types and can be viewed by setting `KMP_AFFINITY=verbose`
@@ -755,21 +609,8 @@ of the OpenMP worker threads at the start of each parallel region.
Selects the OpenMP run-time library execution mode. The values for this variable
are `serial`, `turnaround`, or `throughput`.
-**Default:**
-
-
-
-`throughput`
-
-**Related environment variable:**
-
-
-
-`KMP_BLOCKTIME`
-
- and
-
-`OMP_WAIT_POLICY`
+**Default:** `throughput`\
+**Related environment variable:** `KMP_BLOCKTIME` and `OMP_WAIT_POLICY`
#### KMP_SETTINGS
@@ -813,7 +654,7 @@ Possible values are:
- `cpuid_leaf4` (x86 only) - Decodes the APIC identifiers as specified in leaf 4
of the cpuid instruction. The runtime will produce an error if the machine does not support leaf 4.
- `cpuinfo` - If `KMP_CPUINFO_FILE` is not specified, forces OpenMP to
- parse {file}`/proc/cpuinfo` to determine the topology (Linux only).
+ parse `/proc/cpuinfo` to determine the topology (Linux only).
If `KMP_CPUINFO_FILE` is specified as described above, uses it (Windows or Linux).
- `group` - Models the machine as a 2-level map, with level 0 specifying the
different processors in a group, and level 1 specifying the different
@@ -849,7 +690,7 @@ OpenMP run-time library during program execution.
## LLVM/OpenMP Target Host Runtime (`libomptarget`)
-(libopenmptarget-environment-vars)=
+(libopenmptarget_environment_vars)=
### Environment Variables
@@ -913,7 +754,7 @@ freed after the device kernel exits. The default threshold value is `8KB`. If
manager will be completely disabled.
This has no effect on the host or shared memory managers.
-(libomptarget-info)=
+(libomptarget_info)=
#### LIBOMPTARGET_INFO
@@ -1049,7 +890,7 @@ Info: Removing map entry with HstPtrBegin=0x00007fff0d259a40,
From this information, we can see the OpenMP kernel being launched on the CUDA
device with enough threads and blocks for all `1024` iterations of the loop in
-simplified {doc}`SPMD Mode <Offloading>`. The information from the OpenMP data
+simplified [SPMD Mode](Offloading.md). The information from the OpenMP data
region shows the two arrays `X` and `Y` being copied from the host to the
device. This creates an entry in the host-device mapping table associating the
host pointers to the newly created device data. The data mappings in the OpenMP
@@ -1078,7 +919,7 @@ int main() {
}
```
-(libopenmptarget-errors)=
+(libopenmptarget_errors)=
### Errors:
@@ -1178,12 +1019,12 @@ allocated using `malloc` and `free` for the CUDA plugin. This is necessary
for some applications that allocate too much memory either through the user or
globalization.
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
Offloading
-```
+:::
#### LIBOMPTARGET_MAP_FORCE_ATOMIC
@@ -1215,7 +1056,7 @@ TR14's `attach(always)` map-type-modifier), thereby treating
experimentation, or as a workaround for programs compiled without
`-fopenmp-version=61`.
-(libomptarget-jit-opt-level)=
+(libomptarget_jit_opt_level)=
#### LIBOMPTARGET_JIT_OPT_LEVEL
@@ -1238,7 +1079,7 @@ assembler in object format for the respective target. The JIT optimization
pipeline and backend are skipped and only target specific post-processing is
performed on the object file before it is loaded onto the device.
-(libomptarget-jit-replacement-module)=
+(libomptarget_jit_replacement_module)=
#### LIBOMPTARGET_JIT_REPLACEMENT_MODULE
@@ -1252,7 +1093,7 @@ llvm tools (llvm-objdump), or, simply, by setting the
{ref}`LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE` or
{ref}`LIBOMPTARGET_JIT_POST_OPT_IR_MODULE` environment variables.
-(libomptarget-jit-pre-opt-ir-module)=
+(libomptarget_jit_pre_opt_ir_module)=
#### LIBOMPTARGET_JIT_PRE_OPT_IR_MODULE
@@ -1263,7 +1104,7 @@ which the LLVM-IR module is written. The module can be the analyzed, and
transformed and loaded back into the JIT pipeline via
{ref}`LIBOMPTARGET_JIT_REPLACEMENT_MODULE`.
-(libomptarget-jit-post-opt-ir-module)=
+(libomptarget_jit_post_opt_ir_module)=
#### LIBOMPTARGET_JIT_POST_OPT_IR_MODULE
@@ -1274,7 +1115,7 @@ which the LLVM-IR module is written. The module can be the analyzed, and
transformed and loaded back into the JIT pipeline via
{ref}`LIBOMPTARGET_JIT_REPLACEMENT_MODULE`.
-(libomptarget-jit-save-image-filename)=
+(libomptarget_jit_save_image_filename)=
#### LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME
@@ -1309,7 +1150,7 @@ double-free.
This environment variable determines how many stack traces of kernel launches
are tracked to aid in error reporting, e.g., what asynchronous kernel failed.
-(libomptarget-kernel-record-replay)=
+(libomptarget_kernel_record_replay)=
### Kernel Record Replay
@@ -1398,7 +1239,7 @@ is provided below.
- `LIBOMPTARGET_RECORD_DEVICE=<Num> (default 0)`
- `LIBOMPTARGET_RECORD_OUTPUT=[TRUE/FALSE] (default TRUE)`
-(libomptarget-record)=
+(libomptarget_record)=
#### LIBOMPTARGET_RECORD
@@ -1408,7 +1249,7 @@ overhead to the recorded program. When the recording is disabled, the following
recording environment variables are not considered. The recording is disabled by
default.
-(libomptarget-record-dir)=
+(libomptarget_record_dir)=
#### LIBOMPTARGET_RECORD_DIR
@@ -1416,7 +1257,7 @@ This environment variable is used to specify the relative or absolute path to
the directory where the recorded files will be stored. If omitted or empty, the
files will be stored in current working directory.
-(libomptarget-record-report)=
+(libomptarget_record_report)=
#### LIBOMPTARGET_RECORD_REPORT
@@ -1426,7 +1267,7 @@ report is emitted in the standard output. See
{ref}`LIBOMPTARGET_RECORD_REPORT_FILENAME` to emit the report to a file. By
default, no report is emitted.
-(libomptarget-record-report-filename)=
+(libomptarget_record_report_filename)=
#### LIBOMPTARGET_RECORD_REPORT_FILENAME
@@ -1454,7 +1295,7 @@ recorded. The default device is `0`.
This environment variable is used to instruct the runtime to record the output
device memory snapshot into a file. The default value is `TRUE`.
-(libomptarget-plugin)=
+(libomptarget_plugin)=
## LLVM/OpenMP Target Host Runtime Plugins (`libomptarget.rtl.XXXX`)
@@ -1551,7 +1392,7 @@ queue. If this is disabled, each time a stream is requested a new HSA queue
will be initialized, regardless of their utilization. Additionally, queues will
be selected using round robin selection. The default value is `true`.
-(libomptarget-amdgpu-teams-per-cu)=
+(libomptarget_amdgpu_teams_per_cu)=
#### LIBOMPTARGET_AMDGPU_TEAMS_PER_CU
@@ -1584,7 +1425,7 @@ HSA wait state within the AMDGPU plugin. For the duration of this value
the HSA runtime may busy wait. This can reduce overall latency.
The default value is `2000000`.
-(remote-offloading-plugin)=
+(remote_offloading_plugin)=
### Remote Offloading Plugin:
@@ -1644,7 +1485,7 @@ supported anymore. Please use the `dyn_groupprivate` clause instead, as
shown in {ref}`libomptarget_dynamic_shared`.
:::
-(libomptarget-libc)=
+(libomptarget_libc)=
### LLVM/OpenMP support for C library routines
@@ -1690,7 +1531,7 @@ PluginInterface --> Running an RPC server on device 0
Hello World!
```
-(libomptarget-device)=
+(libomptarget_device)=
## LLVM/OpenMP Target Device Runtime (`libomptarget-ARCH-SUBARCH.bc`)
@@ -1698,7 +1539,7 @@ 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-dynamic-shared)=
+(libomptarget_dynamic_shared)=
### Dynamic Shared Memory
@@ -1734,7 +1575,7 @@ void foo(int N) {
}
```
-(libomptarget-device-allocator)=
+(libomptarget_device_allocator)=
### Device Allocation
@@ -1742,7 +1583,7 @@ The device runtime supports basic runtime allocation via the `omp_alloc`
function. Currently, this allocates global memory for all default traits. Access
modifiers are currently not supported and return a null pointer.
-(libomptarget-device-debugging)=
+(libomptarget_device_debugging)=
### Debugging
@@ -1762,4 +1603,3 @@ debugging features are supported.
[chrome tracing]: https://www.chromium.org/developers/how-tos/trace-event-profiling-tool
[llvm support library]: https://llvm.org/docs/SupportLibrary.html
[speedscope app]: https://www.speedscope.app/
-
diff --git a/openmp/docs/index.md b/openmp/docs/index.md
index 97650a31a36bd9..57b2e92602dfbc 100644
--- a/openmp/docs/index.md
+++ b/openmp/docs/index.md
@@ -8,13 +8,13 @@ additions. Please post on the [Discourse forums (Runtimes -
OpenMP)](https://discourse.llvm.org/c/runtimes/openmp/35)..
:::
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
LLVM/OpenMP Documentation <self>
Building
-```
+:::
# Getting Started
@@ -42,12 +42,12 @@ multitude of available {ref}`OpenMP runtimes <openmp_runtimes>`.
A high-level overview of OpenMP in LLVM can be found {doc}`here <design/Overview>`.
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
design/Overview
-```
+:::
# OpenACC Support
@@ -57,12 +57,12 @@ being extended to serve as OpenACC runtimes. In some cases, Clang
supports {doc}`OpenMP extensions <openacc/OpenMPExtensions>` to make
the additional functionality also available in OpenMP applications.
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
openacc/Overview
-```
+:::
# LLVM/OpenMP Optimizations
@@ -73,35 +73,35 @@ boundaries <optimizations/OpenMPUnawareOptimizations>`.
In-depth discussion of the topic can be found {doc}`here <optimizations/Overview>`.
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
optimizations/Overview
-```
+:::
# LLVM/OpenMP Optimization Remarks
LLVM has an elaborate ecosystem around [analysis and optimization remarks](https://llvm.org/docs/Remarks.html) issues during
-compilation. The remarks can be enabled from the clang frontend [[1]][[1]] [[2]][[2]]
-in various formats [[3]][[3]] [[4]][[4]] to be used by tools, i.a., `opt-viewer` or
-`llvm-opt-report` (dated).
+compilation. The remarks can be enabled from the clang frontend [\[1\]][remarks-1] [\[2\]][remarks-2]
+in various formats [\[3\]][remarks-3] [\[4\]][remarks-4] to be used by tools, i.a., {title-reference}`opt-viewer` or
+{title-reference}`llvm-opt-report` (dated).
The OpenMP optimizations in LLVM have been developed with remark support as a
priority. For a list of OpenMP specific remarks and more information on them,
please refer to {doc}`remarks/OptimizationRemarks`.
-- [[1]][[1]] <https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports>
-- [[2]][[2]] <https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags>
-- [[3]][[3]] <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file>
-- [[4]][[4]] <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record>
+- [\[1\]][remarks-1] <https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports>
+- [\[2\]][remarks-2] <https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags>
+- [\[3\]][remarks-3] <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file>
+- [\[4\]][remarks-4] <https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record>
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
remarks/OptimizationRemarks
-```
+:::
# OpenMP Command-Line Argument Reference
@@ -112,12 +112,12 @@ we also recommend the OpenMP
page that offers a detailed overview of options specific to OpenMP. It also
contains a list of OpenMP offloading related command-line arguments.
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
CommandLineArgumentReference
-```
+:::
# Support, Getting Involved, and Frequently Asked Questions (FAQ)
@@ -128,12 +128,12 @@ the {doc}`Support and FAQ <SupportAndFAQ>` page.
We also encourage everyone interested in OpenMP in LLVM to {doc}`get involved
<SupportAndFAQ>`.
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
SupportAndFAQ
-```
+:::
# Release Notes
@@ -141,15 +141,14 @@ The current (in-progress) release notes can be found {doc}`here <ReleaseNotes>`
release notes for releases, starting with LLVM 12, will be available on [the
Download Page](https://releases.llvm.org/download.html).
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
In-Progress ReleaseNotes <ReleaseNotes>
-```
-
-[[1]]: https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports
-[[2]]: https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags
-[[3]]: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file
-[[4]]: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record
+:::
+[remarks-1]: https://clang.llvm.org/docs/UsersManual.html#options-to-emit-optimization-reports
+[remarks-2]: https://clang.llvm.org/docs/ClangCommandLineReference.html#diagnostic-flags
+[remarks-3]: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-foptimization-record-file
+[remarks-4]: https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang1-fsave-optimization-record
diff --git a/openmp/docs/openacc/OpenMPExtensions.md b/openmp/docs/openacc/OpenMPExtensions.md
index 08516eb05ffff9..bc70160b14ee59 100644
--- a/openmp/docs/openacc/OpenMPExtensions.md
+++ b/openmp/docs/openacc/OpenMPExtensions.md
@@ -35,7 +35,7 @@ OpenMP extensions:
## `ompx_hold` Map Type Modifier
-(ompx-holdexample)=
+(ompx-holdExample)=
### Example
@@ -54,7 +54,7 @@ associated region regardless of any `target exit data` directives
executed during the call to `foo`. Thus, the presence assertion for
`x` at the enclosed `target` construct cannot fail.
-(ompx-holdbehavior)=
+(ompx-holdBehavior)=
### Behavior
@@ -131,7 +131,7 @@ the runtime level. That is, OpenACC's dynamic reference count is
OpenMP's dynamic reference count, and OpenACC's structured reference
count is our OpenMP hold reference count extension.
-(atomicwithinteams)=
+(atomicWithinTeams)=
## `atomic` Strictly Nested Within `teams`
@@ -160,4 +160,3 @@ constructs. For example, the following is conforming OpenACC:
#pragma acc atomic update
x++;
```
-
diff --git a/openmp/docs/openacc/Overview.md b/openmp/docs/openacc/Overview.md
index 2d466428394958..d09538ccf97440 100644
--- a/openmp/docs/openacc/Overview.md
+++ b/openmp/docs/openacc/Overview.md
@@ -4,11 +4,10 @@ OpenACC support is under development for both Flang and Clang. For
this purpose, LLVM's OpenMP runtimes are being extended to serve as
OpenACC runtimes.
-```{toctree}
+:::{toctree}
:glob: true
:hidden: true
:maxdepth: 1
OpenMPExtensions
-```
-
+:::
diff --git a/openmp/docs/optimizations/OpenMPOpt.md b/openmp/docs/optimizations/OpenMPOpt.md
index 8e9924da2138b0..ee72f96e593da1 100644
--- a/openmp/docs/optimizations/OpenMPOpt.md
+++ b/openmp/docs/optimizations/OpenMPOpt.md
@@ -6,14 +6,14 @@ optimization pass will attempt to optimize the module with OpenMP-specific
domain-knowledge. This pass is enabled by default at high optimization levels
(O2 / O3) if compiling with OpenMP support enabled.
-(openmpopt)=
+(OpenMPOpt)=
## OpenMPOpt
-```{contents}
-:depth: 1
+:::{contents}
:local: true
-```
+:depth: 1
+:::
OpenMPOpt contains several OpenMP-Aware optimizations. This pass is run early on
the entire Module, and later on the entire call graph. Most optimizations done
@@ -42,7 +42,8 @@ this it must either be placed in global or shared memory. This needs to be done
every time a variable may potentially be shared in order to create correct
OpenMP programs. Unfortunately, this has significant performance implications
and is not needed in the majority of cases. For example, when Clang is
-generating code for this offloading region, it will see that the variable `x`
+generating code for this offloading region, it will see that the variable
+{title-reference}`x`
escapes and is potentially shared. This will require globalizing the variable,
which means it cannot reside in the registers on the device.
@@ -100,4 +101,3 @@ should be treated as a defect in the program.
- 2021 OpenMP Webinar: "A Compiler's View of OpenMP" <https://youtu.be/eIMpgez61r4>
- 2020 LLVM Developers’ Meeting: "(OpenMP) Parallelism-Aware Optimizations" <https://youtu.be/gtxWkeLCxmU>
- 2019 EuroLLVM Developers’ Meeting: "Compiler Optimizations for (OpenMP) Target Offloading to GPUs" <https://youtu.be/3AbS82C3X30>
-
diff --git a/openmp/docs/optimizations/Overview.md b/openmp/docs/optimizations/Overview.md
index 4559b448704abb..9d6ba301b0275e 100644
--- a/openmp/docs/optimizations/Overview.md
+++ b/openmp/docs/optimizations/Overview.md
@@ -1,13 +1,13 @@
-(llvm-openmp-optimizations)=
+(llvm_openmp_optimizations)=
# OpenMP Optimizations in LLVM
LLVM, since [version 11](https://releases.llvm.org/download.html#11.0.0) (12 Oct
-2020), has an {doc}`OpenMP-Aware optimization pass <OpenMPOpt>`
-as well as the ability to {doc}`perform "scalar optimizations" across OpenMP region
-boundaries <OpenMPUnawareOptimizations>`.
+2020), has an [OpenMP-Aware optimization pass](OpenMPOpt.md)
+as well as the ability to [perform "scalar optimizations" across OpenMP region
+boundaries](OpenMPUnawareOptimizations.md).
-```{toctree}
+:::{toctree}
:glob: true
:hidden: true
:maxdepth: 1
@@ -15,5 +15,4 @@ boundaries <OpenMPUnawareOptimizations>`.
OpenMPOpt
OpenMPUnawareOptimizations
-```
-
+:::
diff --git a/openmp/docs/remarks/OMP100.md b/openmp/docs/remarks/OMP100.md
index ce475d7487be80..0e0ce41614d586 100644
--- a/openmp/docs/remarks/OMP100.md
+++ b/openmp/docs/remarks/OMP100.md
@@ -1,8 +1,8 @@
(omp100)=
-(omp-no-external-caller-in-target-region)=
+(omp_no_external_caller_in_target_region)=
-# Potentially unknown OpenMP target region caller `[OMP100]`
+# Potentially unknown OpenMP target region caller {title-reference}`[OMP100]`
A function remark that indicates the function, when compiled for a GPU, is
potentially called from outside the translation unit. Note that a remark is
@@ -18,10 +18,9 @@ passing scheme and often improve the register usage on the GPU. However, If a
parallel region on the GPU is in a function with external linkage we may not
know all callers statically. If there are outside callers within target
regions, this remark is to be ignored. If there are no such callers, users can
-modify the linkage and thereby help optimization with a `static` or
-`__attribute__((internal))` function annotation. If changing the linkage is
+modify the linkage and thereby help optimization with a {title-reference}`static` or
+{title-reference}`__attribute__((internal))` function annotation. If changing the linkage is
impossible, e.g., because there are outside callers on the host, one can split
the function into an external visible interface which is not compiled for
the target and an internal implementation which is compiled for the target
and should be called from within the target region.
-
diff --git a/openmp/docs/remarks/OMP113.md b/openmp/docs/remarks/OMP113.md
index 9af7c3a3178623..a1870dbbc1de8f 100644
--- a/openmp/docs/remarks/OMP113.md
+++ b/openmp/docs/remarks/OMP113.md
@@ -1,6 +1,6 @@
(omp113)=
-# Could not move globalized variable to the stack. Variable is potentially captured in call. Mark parameter as `__attribute__((noescape))` to override. [OMP113]
+# Could not move globalized variable to the stack. Variable is potentially captured in call. Mark parameter as {title-reference}`__attribute__((noescape))` to override. [OMP113]
This missed remark indicates that a globalized value could not be moved to the
stack because it is potentially captured by a call to a function we cannot
diff --git a/openmp/docs/remarks/OMP121.md b/openmp/docs/remarks/OMP121.md
index 1df315abcd02e0..79d37d38b0d98d 100644
--- a/openmp/docs/remarks/OMP121.md
+++ b/openmp/docs/remarks/OMP121.md
@@ -1,6 +1,6 @@
(omp121)=
-# Value has potential side effects preventing SPMD-mode execution. Add `[[omp::assume(\"ompx_spmd_amenable\")]]` to the called function to override. [OMP121]
+# Value has potential side effects preventing SPMD-mode execution. Add `[[omp::assume("ompx_spmd_amenable")]]` to the called function to override. [OMP121]
This analysis remarks indicates that a potential side-effect that cannot be
guarded prevents the target region from executing in SPMD-mode. SPMD-mode
@@ -74,4 +74,3 @@ omp121.cpp:6:1: remark: Transformed generic-mode kernel to SPMD-mode. [OMP120]
## Diagnostic Scope
OpenMP target offloading analysis remark.
-
diff --git a/openmp/docs/remarks/OMP133.md b/openmp/docs/remarks/OMP133.md
index 59719dff747a68..d175574272f7f7 100644
--- a/openmp/docs/remarks/OMP133.md
+++ b/openmp/docs/remarks/OMP133.md
@@ -1,4 +1,4 @@
-# Call may contain unknown parallel regions. Use `[[omp::assume("omp_no_parallelism")]]` to override. [OMP133]
+# Call may contain unknown parallel regions. Use {title-reference}`[[omp::assume("omp_no_parallelism")]]` to override. [OMP133]
(omp133)=
@@ -65,4 +65,3 @@ omp133.cpp:4:1: remark: Rewriting generic-mode kernel with a customized state ma
## Diagnostic Scope
OpenMP target offloading analysis remark.
-
diff --git a/openmp/docs/remarks/OMP150.md b/openmp/docs/remarks/OMP150.md
index ed8fe9b95d5416..8bec489d0a2134 100644
--- a/openmp/docs/remarks/OMP150.md
+++ b/openmp/docs/remarks/OMP150.md
@@ -37,4 +37,3 @@ omp150.cpp:2:1: remark: Parallel region merged with parallel region at merge.cpp
## Diagnostic Scope
OpenMP optimization remark.
-
diff --git a/openmp/docs/remarks/OMP180.md b/openmp/docs/remarks/OMP180.md
index 36513d1db57de0..9f5fc8077b74e5 100644
--- a/openmp/docs/remarks/OMP180.md
+++ b/openmp/docs/remarks/OMP180.md
@@ -12,7 +12,7 @@ single value after analysis.
This optimization will trigger for most target regions to simplify the runtime
once certain constants are known. This will trigger for internal runtime
functions so it requires enabling verbose remarks with
-`-openmp-opt-verbose-remarks` (prefixed with `-mllvm` for use with clang).
+{title-reference}`-openmp-opt-verbose-remarks` (prefixed with {title-reference}`-mllvm` for use with clang).
```c++
void foo() {
@@ -33,4 +33,3 @@ remark: Replacing runtime call __kmpc_parallel_level with 1. [OMP180] [-Rpass=op
## Diagnostic Scope
OpenMP optimization remark.
-
diff --git a/openmp/docs/remarks/OMP190.md b/openmp/docs/remarks/OMP190.md
index cbcd9ebbeb5d02..9a501cb5d89b69 100644
--- a/openmp/docs/remarks/OMP190.md
+++ b/openmp/docs/remarks/OMP190.md
@@ -13,7 +13,7 @@ Overhead Execution", IPDPS'22.
This optimization will trigger for most target regions if state initialization
was removed as a consequence of "state forwarding". This will trigger for
internal runtime functions so it requires enabling verbose remarks with
-`-openmp-opt-verbose-remarks` (prefixed with `-mllvm` for use with clang).
+{title-reference}`-openmp-opt-verbose-remarks` (prefixed with {title-reference}`-mllvm` for use with clang).
## Diagnostic Scope
diff --git a/openmp/docs/remarks/OptimizationRemarks.md b/openmp/docs/remarks/OptimizationRemarks.md
index eb14376833c1cd..4deea9ab0a6359 100644
--- a/openmp/docs/remarks/OptimizationRemarks.md
+++ b/openmp/docs/remarks/OptimizationRemarks.md
@@ -13,7 +13,7 @@ features of the remark system, consult the clang documentation:
## OpenMP Remarks
-```{toctree}
+:::{toctree}
:hidden: true
:maxdepth: 1
@@ -36,81 +36,79 @@ OMP160
OMP170
OMP180
OMP190
-```
+:::
-```{eval-rst}
-.. list-table::
- :widths: 15 15 70
- :header-rows: 1
+:::{list-table}
+:widths: 15 15 70
+:header-rows: 1
* - Diagnostics Number
- Diagnostics Kind
- Diagnostics Description
- * - :ref:`OMP100 <omp100>`
+ * - {ref}`OMP100 <omp100>`
- Analysis
- Potentially unknown OpenMP target region caller.
- * - :ref:`OMP101 <omp101>`
+ * - {ref}`OMP101 <omp101>`
- Analysis
- Parallel region is used in unknown / unexpected ways. Will not attempt to
rewrite the state machine.
- * - :ref:`OMP102 <omp102>`
+ * - {ref}`OMP102 <omp102>`
- Analysis
- Parallel region is not called from a unique kernel. Will not attempt to
rewrite the state machine.
- * - :ref:`OMP110 <omp110>`
+ * - {ref}`OMP110 <omp110>`
- Optimization
- Moving globalized variable to the stack.
- * - :ref:`OMP111 <omp111>`
+ * - {ref}`OMP111 <omp111>`
- Optimization
- Replaced globalized variable with X bytes of shared memory.
- * - :ref:`OMP112 <omp112>`
+ * - {ref}`OMP112 <omp112>`
- Missed
- Found thread data sharing on the GPU. Expect degraded performance due to
data globalization.
- * - :ref:`OMP113 <omp113>`
+ * - {ref}`OMP113 <omp113>`
- Missed
- Could not move globalized variable to the stack. Variable is potentially
- captured in call. Mark parameter as `__attribute__((noescape))` to
+ captured in call. Mark parameter as {title-reference}`__attribute__((noescape))` to
override.
- * - :ref:`OMP120 <omp120>`
+ * - {ref}`OMP120 <omp120>`
- Optimization
- Transformed generic-mode kernel to SPMD-mode.
- * - :ref:`OMP121 <omp121>`
+ * - {ref}`OMP121 <omp121>`
- Analysis
- Value has potential side effects preventing SPMD-mode execution. Add
- `[[omp::assume(\"ompx_spmd_amenable\")]]` to the called function
+ {title-reference}`[[omp::assume(\"ompx_spmd_amenable\")]]` to the called function
to override.
- * - :ref:`OMP130 <omp130>`
+ * - {ref}`OMP130 <omp130>`
- Optimization
- Removing unused state machine from generic-mode kernel.
- * - :ref:`OMP131 <omp131>`
+ * - {ref}`OMP131 <omp131>`
- Optimization
- Rewriting generic-mode kernel with a customized state machine.
- * - :ref:`OMP132 <omp132>`
+ * - {ref}`OMP132 <omp132>`
- Analysis
- Generic-mode kernel is executed with a customized state machine that
requires a fallback.
- * - :ref:`OMP133 <omp133>`
+ * - {ref}`OMP133 <omp133>`
- Analysis
- Call may contain unknown parallel regions. Use
- `[[omp::assume("omp_no_parallelism")]]` to override.
- * - :ref:`OMP140 <omp140>`
+ {title-reference}`[[omp::assume("omp_no_parallelism")]]` to override.
+ * - {ref}`OMP140 <omp140>`
- Analysis
- Could not internalize function. Some optimizations may not be possible.
- * - :ref:`OMP150 <omp150>`
+ * - {ref}`OMP150 <omp150>`
- Optimization
- - Parallel region merged with parallel region at <location>.
- * - :ref:`OMP160 <omp160>`
+ - Parallel region merged with parallel region at \<location>.
+ * - {ref}`OMP160 <omp160>`
- Optimization
- Removing parallel region with no side-effects.
- * - :ref:`OMP170 <omp170>`
+ * - {ref}`OMP170 <omp170>`
- Optimization
- - OpenMP runtime call <call> deduplicated.
- * - :ref:`OMP180 <omp180>`
+ - OpenMP runtime call \<call> deduplicated.
+ * - {ref}`OMP180 <omp180>`
- Optimization
- - Replacing OpenMP runtime call <call> with <value>.
- * - :ref:`OMP190 <omp190>`
+ - Replacing OpenMP runtime call \<call> with \<value>.
+ * - {ref}`OMP190 <omp190>`
- Optimization
- Redundant barrier eliminated. (device only)
-```
-
+:::
>From ed1b3b55e0f7e01c184d1fb313cf46b3d62c72e3 Mon Sep 17 00:00:00 2001
From: Reid Kleckner <rkleckner at nvidia.com>
Date: Tue, 15 Sep 2026 07:47:04 -0700
Subject: [PATCH 3/3] Apply suggestion from @Meinersbur
Co-authored-by: Michael Kruse <llvm-project at meinersbur.de>
---
openmp/docs/README.txt | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/openmp/docs/README.txt b/openmp/docs/README.txt
index 27d478d001d68c..379143de2552a5 100644
--- a/openmp/docs/README.txt
+++ b/openmp/docs/README.txt
@@ -16,7 +16,7 @@ Sphinx <http://sphinx-doc.org/> and then do:
make docs-openmp-html
$BROWSER <build-dir>/docs/html/index.html
-The mapping between reStructuredText files and generated documentation is
+The mapping between Markdown files and generated documentation is
`docs/Foo.md` <-> `<build-dir>/projects/openmp/docs//html/Foo.html` <->
`https://openmp.llvm.org/docs/Foo.html`.
More information about the llvm-branch-commits
mailing list