[OpenACC][OpenMP] Document atomic-in-teams extension

[jdenny]

Summary

That is, put D126323 in the status doc and explain its relationship to
OpenACC support.

[jdenny]

Thanks for the quick review.

clang/docs/OpenMPSupport.rst

 
 The following table provides a quick overview over various OpenMP
 extensions and their implementation status.  These extensions are not
 currently defined by any standard, so links to associated LLVM
 documentation are provided.  As these extensions mature, they will be
 considered for standardization.  Please contact *openmp-dev* at
 *lists.llvm.org* to provide feedback.
 
-+------------------------------+---------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
-|Category                      | Feature                                                                   | Status                   | Reviews                                                |
-+==============================+===========================================================================+==========================+========================================================+
-| device extension             | `'ompx_hold' map type modifier                                            | :good:`prototyped`       | D106509, D106510                                       |
-|                              | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#ompx-hold>`_  |                          |                                                        |
-+------------------------------+---------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
++------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
+|Category                      | Feature                                                                           | Status                   | Reviews                                                |
++==============================+===================================================================================+==========================+========================================================+
+| atomic extension             | `'atomic' strictly nested within 'teams'                                          | :good:`prototyped`       | D126323                                                |
+|                              | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#atomicWithinTeams>`_  |                          |                                                        |
++------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
+| device extension             | `'ompx_hold' map type modifier                                                    | :good:`prototyped`       | D106509, D106510                                       |
+|                              | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#ompx-hold>`_          |                          |                                                        |
++------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+

openmp/docs/openacc/OpenMPExtensions.rst

 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++;