[OpenMP] Extend omp teams to permit nested omp tile

[jdenny]

Summary

OpenMP 5.2, sec. 10.2 "teams Construct", p. 232, L9-12 restricts what
regions can be strictly nested within a teams construct. This patch
relaxes Clang's enforcement of this restriction in the case of nested
tile constructs unless -fno-openmp-extensions is specified. Cases
like the following then seem to work fine with no additional
implementation changes:

#pragma omp target teams
#pragma omp tile sizes(N, M)
for (int i = 0; i < I; ++i) {
  for (int j = 0; j < J; ++j) {
    ...
  }
}

This commit is similar to D126323 (48ca3a5ebb15) plus D126547
(4a368136693b), which relaxed the restriction for an atomic
construct in a teams construct.

[jdenny]

Added phab review number to status doc.

[jdenny]

Ping.

clang/docs/OpenMPSupport.rst

 documentation are provided.  As these extensions mature, they will be
 considered for standardization. Please post on the
 `Discourse forums (Runtimes - OpenMP category)`_ to provide feedback.
 
 +------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
 |Category                      | Feature                                                                           | Status                   | Reviews                                                |
 +==============================+===================================================================================+==========================+========================================================+
 | atomic extension             | `'atomic' strictly nested within 'teams'                                          | :good:`prototyped`       | D126323                                                |
-|                              | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#atomicWithinTeams>`_  |                          |                                                        |
+|                              | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#withinTeams>`_        |                          |                                                        |
 +------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
 | device extension             | `'ompx_hold' map type modifier                                                    | :good:`prototyped`       | D106509, D106510                                       |
 |                              | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#ompx-hold>`_          |                          |                                                        |
 +------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
+| tile extension               | `'tile' strictly nested within 'teams'                                            | :good:`prototyped`       | D151350                                                |
+|                              | <https://openmp.llvm.org/docs/openacc/OpenMPExtensions.html#withinTeams>`_        |                          |                                                        |
++------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+
 
 .. _Discourse forums (Runtimes - OpenMP category): https://discourse.llvm.org/c/runtimes/openmp/35

clang/lib/Sema/SemaOpenMP.cpp

if (!NestingProhibited &&
       // OpenMP [5.1, 2.22, Nesting of Regions]
       // distribute, distribute simd, distribute parallel worksharing-loop,
       // distribute parallel worksharing-loop SIMD, loop, parallel regions,
       // including any parallel regions arising from combined constructs,
       // omp_get_num_teams() regions, and omp_get_team_num() regions are the
       // only OpenMP regions that may be strictly nested inside the teams
       // region.
       //
-      // As an extension, we permit atomic within teams as well.
-      NestingProhibited = !isOpenMPParallelDirective(CurrentRegion) &&
-                          !isOpenMPDistributeDirective(CurrentRegion) &&
-                          CurrentRegion != OMPD_loop &&
-                          !(SemaRef.getLangOpts().OpenMPExtensions &&
-                            CurrentRegion == OMPD_atomic);
+      // As an extension, we permit atomic and tile within teams as well.
+      NestingProhibited =
+          !isOpenMPParallelDirective(CurrentRegion) &&
+            (CurrentRegion == OMPD_atomic || CurrentRegion == OMPD_tile));
       Recommend = ShouldBeInParallelRegion;
     }
     if (!NestingProhibited && CurrentRegion == OMPD_loop) {
       // OpenMP [5.1, 2.11.7, loop Construct, Restrictions]
       // If the bind clause is present on the loop construct and binding is
       // teams then the corresponding loop region must be strictly nested inside
       // a teams region.
       NestingProhibited = BindKind == OMPC_BIND_teams &&

openmp/docs/openacc/OpenMPExtensions.rst

 
 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:
+.. _withinTeams:
 
-``atomic`` Strictly Nested Within ``teams``
--------------------------------------------
+Regions 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
+``atomic`` or ``tile`` construct.  For example, the following case is
 permitted:
 
 .. code-block:: c++
 
   #pragma omp target teams map(tofrom:x)
   #pragma omp atomic update
   x++;
 

openmp/libomptarget/test/offloading/target-teams-tile.c

+// Check that omp tile (introduced in OpenMP 5.1) is permitted and behaves when
+// strictly nested within omp target teams.  This is an extension to OpenMP 5.2
+// and is enabled by default.
+
+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+#define NUM_TEAMS_UPPER 256
+#define I_NTILES 8
+#define J_NTILES 9
+#define I_NELEMS 2
+#define J_NELEMS 3
+
+int main() {
+  int numTeams;
+  int order[NUM_TEAMS_UPPER][I_NTILES][J_NTILES][I_NELEMS][J_NELEMS];
+  #pragma omp target teams num_teams(NUM_TEAMS_UPPER) map(from : numTeams)
+  {
+    int team = omp_get_team_num();
+    if (team == 0)
+      numTeams = omp_get_num_teams();
+    int next = 0;
+    #pragma omp tile sizes(I_NELEMS, J_NELEMS)
+    for (int i = 0; i < I_NTILES * I_NELEMS; ++i) {
+      for (int j = 0; j < J_NTILES * J_NELEMS; ++j) {
+        int iTile = i / I_NELEMS;
+        int jTile = j / J_NELEMS;
+        int iElem = i % I_NELEMS;
+        int jElem = j % J_NELEMS;
+        order[team][iTile][jTile][iElem][jElem] = next++;
+      }
+    }
+  }
+  printf("numTeams = %d\n", numTeams);
+  for (int team = 0; team < numTeams; ++team) {
+    int expected = 0;
+    for (int iTile = 0; iTile < I_NTILES; ++iTile) {
+      for (int jTile = 0; jTile < J_NTILES; ++jTile) {
+        for (int iElem = 0; iElem < I_NELEMS; ++iElem) {
+          for (int jElem = 0; jElem < J_NELEMS; ++jElem) {
+            int actual = order[team][iTile][jTile][iElem][jElem];
+            if (expected != actual) {
+              printf("error: order[%d][%d][%d][%d][%d] = %d, expected %d\n",
+                     team, iTile, jTile, iElem, jElem, actual, expected);
+              return 1;
+            }
+            ++expected;
+          }
+        }
+      }
+    }
+  }
+  // CHECK: success
+  printf("success\n");
+  return 0;
+}

openmp/runtime/test/teams/teams-tile.c

+// Check that omp tile (introduced in OpenMP 5.1) is permitted and behaves when
+// strictly nested within omp teams.  This is an extension to OpenMP 5.2 and is
+// enabled by default.
+
+// RUN: %libomp-compile -fopenmp-version=51
+// RUN: %libomp-run 2>&1 | FileCheck %s
+
+#include <omp.h>
+#include <stdio.h>
+
+#define NUM_TEAMS_UPPER 256
+#define I_NTILES 8
+#define J_NTILES 9
+#define I_NELEMS 2
+#define J_NELEMS 3
+
+int main() {
+  int numTeams;
+  int order[NUM_TEAMS_UPPER][I_NTILES][J_NTILES][I_NELEMS][J_NELEMS];
+  #pragma omp teams num_teams(NUM_TEAMS_UPPER)
+  {
+    int team = omp_get_team_num();
+    if (team == 0)
+      numTeams = omp_get_num_teams();
+    int next = 0;
+    #pragma omp tile sizes(I_NELEMS, J_NELEMS)
+    for (int i = 0; i < I_NTILES * I_NELEMS; ++i) {
+      for (int j = 0; j < J_NTILES * J_NELEMS; ++j) {
+        int iTile = i / I_NELEMS;
+        int jTile = j / J_NELEMS;
+        int iElem = i % I_NELEMS;
+        int jElem = j % J_NELEMS;
+        order[team][iTile][jTile][iElem][jElem] = next++;
+      }
+    }
+  }
+  printf("numTeams = %d\n", numTeams);
+  for (int team = 0; team < numTeams; ++team) {
+    int expected = 0;
+    for (int iTile = 0; iTile < I_NTILES; ++iTile) {
+      for (int jTile = 0; jTile < J_NTILES; ++jTile) {
+        for (int iElem = 0; iElem < I_NELEMS; ++iElem) {
+          for (int jElem = 0; jElem < J_NELEMS; ++jElem) {
+            int actual = order[team][iTile][jTile][iElem][jElem];
+            if (expected != actual) {
+              printf("error: order[%d][%d][%d][%d][%d] = %d, expected %d\n",
+                     team, iTile, jTile, iElem, jElem, actual, expected);
+              return 1;
+            }
+            ++expected;
+          }
+        }
+      }
+    }
+  }
+  // CHECK: success
+  printf("success\n");
+  return 0;
+}