openmp/docs/openacc/OpenMPExtensions.rst

OpenMP Extensions for OpenACC
=============================

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.

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:

Motivation
----------

There are multiple benefits to exposing OpenACC functionality as LLVM
OpenMP extensions:

* OpenMP applications can take advantage of the additional
  functionality.
* 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
  is otherwise impossible.
* 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
  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`` Map Type Modifier
-------------------------------

.. _ompx-holdExample:

Example
^^^^^^^

.. code-block:: 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);
  }

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:

Behavior
^^^^^^^^

* 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
  data must also reach zero before the data is unmapped.
* 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
  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
  is in effect.
* 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
  modifier.
* ``ompx_hold`` support is not yet provided for ``defaultmap``.

Implementation
^^^^^^^^^^^^^^

* 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
  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
  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
  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
  hold reference count is greater than zero.

Relationship with OpenACC
^^^^^^^^^^^^^^^^^^^^^^^^^

OpenACC specifies two reference counts for tracking host/device data
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.

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
<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
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
OpenMP's dynamic reference count, and OpenACC's structured reference
count is our OpenMP hold reference count extension.

.. _atomicWithinTeams:

``atomic`` Strictly Nested Within ``teams``
-------------------------------------------

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
extension, Clang relaxes that restriction in the case of the
``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++;

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:

.. code-block:: c++

  #pragma acc parallel copy(x)
  #pragma acc atomic update
  x++;