Index: clang/docs/OpenMPSupport.rst =================================================================== --- clang/docs/OpenMPSupport.rst +++ clang/docs/OpenMPSupport.rst @@ -378,10 +378,13 @@ |Category | Feature | Status | Reviews | +==============================+===================================================================================+==========================+========================================================+ | atomic extension | `'atomic' strictly nested within 'teams' | :good:`prototyped` | D126323 | -| | `_ | | | +| | `_ | | | +------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ | device extension | `'ompx_hold' map type modifier | :good:`prototyped` | D106509, D106510 | | | `_ | | | +------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ +| tile extension | `'tile' strictly nested within 'teams' | :good:`prototyped` | D151350 | +| | `_ | | | ++------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ .. _Discourse forums (Runtimes - OpenMP category): https://discourse.llvm.org/c/runtimes/openmp/35 Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -5199,12 +5199,13 @@ // 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) && + !isOpenMPDistributeDirective(CurrentRegion) && + CurrentRegion != OMPD_loop && + !(SemaRef.getLangOpts().OpenMPExtensions && + (CurrentRegion == OMPD_atomic || CurrentRegion == OMPD_tile)); Recommend = ShouldBeInParallelRegion; } if (!NestingProhibited && CurrentRegion == OMPD_loop) { Index: openmp/docs/openacc/OpenMPExtensions.rst =================================================================== --- openmp/docs/openacc/OpenMPExtensions.rst +++ openmp/docs/openacc/OpenMPExtensions.rst @@ -138,10 +138,10 @@ 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 ^^^^^^^ @@ -149,7 +149,7 @@ 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++ Index: openmp/libomptarget/test/offloading/target-teams-tile.c =================================================================== --- /dev/null +++ openmp/libomptarget/test/offloading/target-teams-tile.c @@ -0,0 +1,59 @@ +// 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 +#include + +#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; +} Index: openmp/runtime/test/teams/teams-tile.c =================================================================== --- /dev/null +++ openmp/runtime/test/teams/teams-tile.c @@ -0,0 +1,59 @@ +// 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 +#include + +#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; +}