caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[OpenMP] Extend omp teams to permit nested omp atomic

Browse Source 
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
`atomic` 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 map(tofrom:x)
 #pragma omp atomic update
 x++;
```

Reviewed By: ABataev

Differential Revision: https://reviews.llvm.org/D126323
This commit is contained in:
Joel E. Denny
2022-05-26 14:59:16 -04:00
parent 3bba72e653
commit 48ca3a5ebb
4 changed files with 115 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
clang/lib/Sema/SemaOpenMP.cpp
View File

@@ -4975,9 +4975,13 @@ static bool checkNestingOfRegions(Sema &SemaRef, const DSAStackTy *Stack,
// 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;
CurrentRegion != OMPD_loop &&
!(SemaRef.getLangOpts().OpenMPExtensions &&
CurrentRegion == OMPD_atomic);
Recommend = ShouldBeInParallelRegion;
}
if (!NestingProhibited && CurrentRegion == OMPD_loop) {

21
clang/test/OpenMP/nesting_of_regions.cpp
View File

@@ -1,10 +1,11 @@
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45,omp45warn %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -verify=expected,omp50 %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45 -Wno-openmp %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45 -Wno-source-uses-openmp %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -fno-openmp-extensions -verify=expected,omp45,omp45warn,omp %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fno-openmp-extensions -verify=expected,omp50,omp %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-extensions -verify=expected,omp50 %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45,omp -fno-openmp-extensions -Wno-openmp %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp -fopenmp-version=45 -verify=expected,omp45,omp -fno-openmp-extensions -Wno-source-uses-openmp %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -fopenmp-version=45 -verify=expected,omp45,omp45warn %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -verify=expected,omp50 %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -fopenmp-version=45 -fno-openmp-extensions -verify=expected,omp45,omp45warn,omp %s
// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -verify=expected,omp50,omp -fno-openmp-extensions %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
void bar();
@@ -5264,7 +5265,7 @@ void foo() {
#pragma omp target
#pragma omp teams
{
#pragma omp atomic // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
#pragma omp atomic // omp-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
++a;
}
#pragma omp target
@@ -8422,7 +8423,7 @@ void foo() {
}
#pragma omp target teams
{
#pragma omp atomic // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
#pragma omp atomic // omp-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
++a;
}
#pragma omp target teams
@@ -14096,7 +14097,7 @@ void foo() {
#pragma omp target
#pragma omp teams
{
#pragma omp atomic // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
#pragma omp atomic // omp-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
++a;
}
#pragma omp target
@@ -17299,7 +17300,7 @@ void foo() {
}
#pragma omp target teams
{
#pragma omp atomic // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
#pragma omp atomic // omp-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
++a;
}
#pragma omp target teams

50
openmp/libomptarget/test/offloading/target-teams-atomic.c Normal file
View File

@@ -0,0 +1,50 @@
// Check that omp atomic 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-run-and-check-generic
#include <omp.h>
#include <stdbool.h>
#include <stdio.h>
#include <string.h>
// High parallelism increases our chances of detecting a lack of atomicity.
#define NUM_TEAMS_TRY 256
int main() {
// CHECK: update: num_teams=[[#NUM_TEAMS:]]{{$}}
// CHECK-NEXT: update: x=[[#NUM_TEAMS]]{{$}}
int x = 0;
int numTeams;
#pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom:x, numTeams)
{
#pragma omp atomic update
++x;
if (omp_get_team_num() == 0)
numTeams = omp_get_num_teams();
}
printf("update: num_teams=%d\n", numTeams);
printf("update: x=%d\n", x);
// CHECK-NEXT: capture: x=[[#NUM_TEAMS]]{{$}}
// CHECK-NEXT: capture: xCapturedCount=[[#NUM_TEAMS]]{{$}}
bool xCaptured[numTeams];
memset(xCaptured, 0, sizeof xCaptured);
x = 0;
#pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom:x, numTeams)
{
int v;
#pragma omp atomic capture
v = x++;
xCaptured[v] = true;
}
printf("capture: x=%d\n", x);
int xCapturedCount = 0;
for (int i = 0; i < numTeams; ++i) {
if (xCaptured[i])
++xCapturedCount;
}
printf("capture: xCapturedCount=%d\n", xCapturedCount);
return 0;
}

49
openmp/runtime/test/teams/teams-atomic.c Normal file
View File

@@ -0,0 +1,49 @@
// Check that omp atomic 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-and-run | FileCheck %s
#include <omp.h>
#include <stdbool.h>
#include <stdio.h>
#include <string.h>
// High parallelism increases our chances of detecting a lack of atomicity.
#define NUM_TEAMS_TRY 256
int main() {
// CHECK: update: num_teams=[[#NUM_TEAMS:]]{{$}}
// CHECK-NEXT: update: x=[[#NUM_TEAMS]]{{$}}
int x = 0;
int numTeams;
#pragma omp teams num_teams(NUM_TEAMS_TRY)
{
#pragma omp atomic update
++x;
if (omp_get_team_num() == 0)
numTeams = omp_get_num_teams();
}
printf("update: num_teams=%d\n", numTeams);
printf("update: x=%d\n", x);
// CHECK-NEXT: capture: x=[[#NUM_TEAMS]]{{$}}
// CHECK-NEXT: capture: xCapturedCount=[[#NUM_TEAMS]]{{$}}
bool xCaptured[numTeams];
memset(xCaptured, 0, sizeof xCaptured);
x = 0;
#pragma omp teams num_teams(NUM_TEAMS_TRY)
{
int v;
#pragma omp atomic capture
v = x++;
xCaptured[v] = true;
}
printf("capture: x=%d\n", x);
int xCapturedCount = 0;
for (int i = 0; i < numTeams; ++i) {
if (xCaptured[i])
++xCapturedCount;
}
printf("capture: xCapturedCount=%d\n", xCapturedCount);
return 0;
}